public inbox for gcc-patches@gcc.gnu.org
 help / color / mirror / Atom feed
* [PATCH] Fix ICE for static vars in offloaded functions
@ 2018-03-07 13:20 Tom de Vries
  2018-03-07 22:13 ` Jakub Jelinek
  0 siblings, 1 reply; 7+ messages in thread
From: Tom de Vries @ 2018-03-07 13:20 UTC (permalink / raw)
  To: Jakub Jelinek, Richard Biener; +Cc: GCC Patches

[-- Attachment #1: Type: text/plain, Size: 1919 bytes --]

Hi,

if we compile the testcase pr84592-2.c from the patch:
...
#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;
}
...

for nvptx offloading, we run into an assert:
...
lto1: internal compiler error: in input_varpool_node, at lto-cgraph.c:1424
0x959ebb input_varpool_node
         gcc/lto-cgraph.c:1422
0x959ebb input_cgraph_1
         gcc/lto-cgraph.c:1544
0x959ebb input_symtab()
         gcc/lto-cgraph.c:1858
0x5aceac read_cgraph_and_symbols
         gcc/lto/lto.c:2891
0x5aceac lto_main()
         gcc/lto/lto.c:3356
...

The assert we run into is:
...
1422      gcc_assert (flag_ltrans
1423                  || (!node->in_other_partition
1424                      && !node->used_from_other_partition));
...

where node is:
...
(gdb) call debug_generic_expr (node.decl)
test
...

and the reason the assert triggers is:
...
(gdb) p node.in_other_partition
$1 = 1
...

AFAIU, what this means is that the variable test is placed in a 
different partition than the offloading function main._omp_fn.0 that 
uses the variable.


I looked at where global variables are put into offload_vars, and found 
that that happens in varpool_node::get_create:
...
   if ((flag_openacc || flag_openmp)
       && lookup_attribute ("omp declare target",
                            DECL_ATTRIBUTES (decl)))
     {
       node->offloadable = 1;
       if (ENABLE_OFFLOADING && !DECL_EXTERNAL (decl))
         {
           g->have_offload = true;
           if (!in_lto_p)
             vec_safe_push (offload_vars, decl);
         }
     }

...

The patch fixes the ICE there by marking the varpool_node test as 
offloadable as well.

Build and reg-tested libgomp on x86_64 with nvptx accelerator.
Bootstrapped and reg-tested on x86_64.

OK for stage4 trunk?

Thanks,
- Tom

[-- Attachment #2: 0001-Fix-ICE-for-static-vars-in-offloaded-functions.patch --]
[-- Type: text/x-patch, Size: 3559 bytes --]

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

^ permalink raw reply	[flat|nested] 7+ messages in thread
end of thread, other threads:[~2018-03-26 10:30 UTC | newest]

Thread overview: 7+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2018-03-07 13:20 [PATCH] Fix ICE for static vars in offloaded functions Tom de Vries
2018-03-07 22:13 ` Jakub Jelinek
2018-03-07 13:29   ` Richard Biener
2018-03-07 14:13     ` Jakub Jelinek
2018-03-07 14:51     ` Tom de Vries
2018-03-07 15:01       ` Richard Biener
2018-03-26 11:04         ` Tom de Vries

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).