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
Fix ICE for static vars in offloaded functions
2018-03-06 Tom de Vries <t...@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;
+}
