[PATCH] Fix ICE for static vars in offloaded functions

[PATCH] Fix ICE for static vars in offloaded functions

[Tom de Vries]

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

Re: [PATCH] Fix ICE for static vars in offloaded functions

[Richard Biener]

On Wed, 7 Mar 2018, Jakub Jelinek wrote:

> On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote:
> > 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.
> 
> Ok, thanks

+      bool in_offload_func
+       = (cfun
+          && TREE_STATIC (decl)
+          && (lookup_attribute ("omp target entr

I think you want to use decl_function_context (decl) here, 
not rely on magic cfun being set.  The whole varpool.c file
doesn't mention cfun yet and you shoudln't either.

please fix if you already committed the fix.

Thanks,
Richard.

> 	Jakub
> 
> 

-- 
Richard Biener <rguenther@suse.de>
SUSE LINUX GmbH, GF: Felix Imendoerffer, Jane Smithard, Graham Norton, HRB 21284 (AG Nuernberg)

Re: [PATCH] Fix ICE for static vars in offloaded functions

[Jakub Jelinek]

On Wed, Mar 07, 2018 at 02:29:48PM +0100, Richard Biener wrote:
> On Wed, 7 Mar 2018, Jakub Jelinek wrote:
> 
> > On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote:
> > > 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.
> > 
> > Ok, thanks
> 
> +      bool in_offload_func
> +       = (cfun
> +          && TREE_STATIC (decl)
> +          && (lookup_attribute ("omp target entr
> 
> I think you want to use decl_function_context (decl) here, 
> not rely on magic cfun being set.  The whole varpool.c file
> doesn't mention cfun yet and you shoudln't either.

Oops, sure, thanks for catching it.

	Jakub

Re: [PATCH] Fix ICE for static vars in offloaded functions

[Tom de Vries]

On 03/07/2018 02:29 PM, Richard Biener wrote:
> On Wed, 7 Mar 2018, Jakub Jelinek wrote:
> 
>> On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote:
>>> 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.
>>
>> Ok, thanks
> 
> +      bool in_offload_func
> +       = (cfun
> +          && TREE_STATIC (decl)
> +          && (lookup_attribute ("omp target entr
> 
> I think you want to use decl_function_context (decl) here,
> not rely on magic cfun being set.  The whole varpool.c file
> doesn't mention cfun yet and you shoudln't either.
> 

decl_function_context (decl) returns main:
...
(gdb) call debug_generic_expr (decl)
test
(gdb) call  decl_function_context (decl)
$2 = (tree_node *) 0x7ffff6978c00
(gdb) call debug_generic_expr ($2)
main
...
while the function annotated as being an offload function is main._omp_fn.0.

The varpool_node::get_create is called during cgraph_edge::rebuild_edges 
here in expand_omp_target:
...
7087          /* Fix the callgraph edges for child_cfun.  Those for cfun 
will be
7088             fixed in a following pass.  */
7089          push_cfun (child_cfun);
7090          if (need_asm)
7091            assign_assembler_name_if_needed (child_fn);
7092          cgraph_edge::rebuild_edges ();
...

Thanks,
- Tom

Re: [PATCH] Fix ICE for static vars in offloaded functions

[Richard Biener]

On Wed, 7 Mar 2018, Tom de Vries wrote:

> On 03/07/2018 02:29 PM, Richard Biener wrote:
> > On Wed, 7 Mar 2018, Jakub Jelinek wrote:
> > 
> > > On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote:
> > > > 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.
> > > 
> > > Ok, thanks
> > 
> > +      bool in_offload_func
> > +       = (cfun
> > +          && TREE_STATIC (decl)
> > +          && (lookup_attribute ("omp target entr
> > 
> > I think you want to use decl_function_context (decl) here,
> > not rely on magic cfun being set.  The whole varpool.c file
> > doesn't mention cfun yet and you shoudln't either.
> > 
> 
> decl_function_context (decl) returns main:
> ...
> (gdb) call debug_generic_expr (decl)
> test
> (gdb) call  decl_function_context (decl)
> $2 = (tree_node *) 0x7ffff6978c00
> (gdb) call debug_generic_expr ($2)
> main
> ...
> while the function annotated as being an offload function is main._omp_fn.0.

Well, that's because the static isn't duplicated (it can't be) so it 
retains the original context.

> The varpool_node::get_create is called during cgraph_edge::rebuild_edges here
> in expand_omp_target:

But at this point it's not created but just looked up, right?

I think the fix is to mark the decl as offloaded when we walk the IL
of the outlined function.  The current point looks like a hack.

Richard.

> ...
> 7087          /* Fix the callgraph edges for child_cfun.  Those for cfun will
> be
> 7088             fixed in a following pass.  */
> 7089          push_cfun (child_cfun);
> 7090          if (need_asm)
> 7091            assign_assembler_name_if_needed (child_fn);
> 7092          cgraph_edge::rebuild_edges ();
> ...
> 
> Thanks,
> - Tom
> 
> 

-- 
Richard Biener <rguenther@suse.de>
SUSE LINUX GmbH, GF: Felix Imendoerffer, Jane Smithard, Graham Norton, HRB 21284 (AG Nuernberg)

Re: [PATCH] Fix ICE for static vars in offloaded functions

[Jakub Jelinek]

On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote:
> 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.

Ok, thanks

	Jakub

Re: [PATCH] Fix ICE for static vars in offloaded functions

[Tom de Vries]

On 03/07/2018 04:01 PM, Richard Biener wrote:
> On Wed, 7 Mar 2018, Tom de Vries wrote:
> 
>> On 03/07/2018 02:29 PM, Richard Biener wrote:
>>> On Wed, 7 Mar 2018, Jakub Jelinek wrote:
>>>
>>>> On Wed, Mar 07, 2018 at 02:20:26PM +0100, Tom de Vries wrote:
>>>>> 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.
>>>>
>>>> Ok, thanks
>>>
>>> +      bool in_offload_func
>>> +       = (cfun
>>> +          && TREE_STATIC (decl)
>>> +          && (lookup_attribute ("omp target entr
>>>
>>> I think you want to use decl_function_context (decl) here,
>>> not rely on magic cfun being set.  The whole varpool.c file
>>> doesn't mention cfun yet and you shoudln't either.
>>>
>>
>> decl_function_context (decl) returns main:
>> ...
>> (gdb) call debug_generic_expr (decl)
>> test
>> (gdb) call  decl_function_context (decl)
>> $2 = (tree_node *) 0x7ffff6978c00
>> (gdb) call debug_generic_expr ($2)
>> main
>> ...
>> while the function annotated as being an offload function is main._omp_fn.0.
> 
> Well, that's because the static isn't duplicated (it can't be) so it
> retains the original context.
> 

[ Actually the static is duplicated in replace_by_duplicate_decl, but 
the statements using it are not rewritten to use the duplicate, so 
indeed, effectively it's not duplicated. ]

>> The varpool_node::get_create is called during cgraph_edge::rebuild_edges here
>> in expand_omp_target:
> 
> But at this point it's not created but just looked up, right?
> 

No, the varpool_node is created at that point.

> I think the fix is to mark the decl as offloaded when we walk the IL
> of the outlined function.  The current point looks like a hack.
> 

OK, I'll try to find a better fix location.

Thanks,
- Tom

> Richard.
> 
>> ...
>> 7087          /* Fix the callgraph edges for child_cfun.  Those for cfun will
>> be
>> 7088             fixed in a following pass.  */
>> 7089          push_cfun (child_cfun);
>> 7090          if (need_asm)
>> 7091            assign_assembler_name_if_needed (child_fn);
>> 7092          cgraph_edge::rebuild_edges ();
>> ...
>>
>> Thanks,
>> - Tom
>>
>>
>