Bug 84592 - [openacc,openmp] lto1: ICE in input_varpool_node, at lto-cgraph.c:1424: for function static var
Summary: [openacc,openmp] lto1: ICE in input_varpool_node, at lto-cgraph.c:1424: for f...
Status: RESOLVED DUPLICATE of bug 90779
Alias: None
Product: gcc
Classification: Unclassified
Component: lto (show other bugs)
Version: 8.0
: P3 normal
Target Milestone: ---
Assignee: Not yet assigned to anyone
URL:
Keywords: lto, openacc, openmp, patch
Depends on:
Blocks:
 
Reported: 2018-02-27 14:46 UTC by Tom de Vries
Modified: 2019-06-15 11:51 UTC (History)
2 users (show)

See Also:
Host:
Target: nvptx
Build:
Known to work:
Known to fail:
Last reconfirmed: 2018-02-27 00:00:00


Attachments
Tentative patch (450 bytes, patch)
2018-03-05 17:56 UTC, Tom de Vries
Details | Diff

Note You need to log in before you can comment on or make changes to this bug.
Description Tom de Vries 2018-02-27 14:46:47 UTC
Consider the following openacc testcase:
...
/* { dg-additional-options "-ftree-switch-conversion" } */

#include <stdio.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[1] = 4;
#pragma acc parallel copy(n)
  {
    n[0] = foo (n[0]); 
  }
  printf ("n: %d\n", n[0]);

  return 0;
}
...

When compiling this (by copying it into libgomp/testsuite/libgomp.oacc-c-c++-common), we run into:
...
lto1: internal compiler error: in input_varpool_node, at lto-cgraph.c:1424^M
0x959ebb input_varpool_node^M
        gcc/lto-cgraph.c:1422^M
0x959ebb input_cgraph_1^M
        gcc/lto-cgraph.c:1544^M
0x959ebb input_symtab()^M
        gcc/lto-cgraph.c:1858^M
0x5aceac read_cgraph_and_symbols^M
        gcc/lto/lto.c:2891^M
0x5aceac lto_main()^M
        gcc/lto/lto.c:3356^M
...

When running lto1 using gdb we find the node we're processing is:
...
(gdb) call debug_generic_expr ( node.decl )
CSWTCH.4
...

The CSWTCH.4 is a symbol introduced by -ftree-switch-conversion. If we use -fno-tree-switch-conversion there's no ICE.
Comment 1 Martin Liška 2018-02-27 15:05:31 UTC
Can you please help me how to run the single test on my machine w/o nvptx target?
Comment 2 Martin Liška 2018-02-27 15:08:13 UTC
The problematic symbol looks as follows:

CSWTCH.5/18 (CSWTCH.5) @0x7ffff6786a00
  Type: variable definition analyzed
  Visibility: prevailing_def_ironly asm_written artificial
  Aux: @0x1
  References: 
  Referring: 
  Availability: available
  Varpool flags: initialized used-by-single-function read-only const-value-known
Comment 3 Tom de Vries 2018-02-28 12:36:09 UTC
(In reply to Martin Liška from comment #1)
> Can you please help me how to run the single test on my machine w/o nvptx
> target?

This is the openmp variant:
...
/* { dg-additional-options "-ftree-switch-conversion" } */

#include <stdio.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;
  }
  
  printf ("n: %d\n", n[0]);

  return 0;
}
...

This might also reproduce with intelmic.
Comment 4 Martin Liška 2018-02-28 13:54:44 UTC
Good, now I can reproduce that with intelmic offloading compiler.
Comment 5 Martin Liška 2018-02-28 14:36:18 UTC
So it's equivalent to:

$ cat ~/Programming/testcases/pr84592-3.c
#include <stdio.h>

int
main (void)
{
  int n[1];
  n[0] = 4;
#pragma omp target 
  {
    static const int test[] = {1,2,3,4};
    n[1] += test[n[0]];
  }
  
  printf ("n: %d\n", n[0]);

  return 0;
}

$ /tmp/install/usr/local/bin/x86_64-pc-linux-gnu-gcc ~/Programming/testcases/pr84592-3.c -fopenmp 
setting in_other_partition: test.2531: 1
test: 1
lto1: internal compiler error: in input_varpool_node, at lto-cgraph.c:1427
0x9840dc input_varpool_node
	/home/marxin/Programming/gcc2/gcc/lto-cgraph.c:1425
0x9840dc input_cgraph_1
	/home/marxin/Programming/gcc2/gcc/lto-cgraph.c:1547
0x9840dc input_symtab()
	/home/marxin/Programming/gcc2/gcc/lto-cgraph.c:1861
0x609f90 read_cgraph_and_symbols
	/home/marxin/Programming/gcc2/gcc/lto/lto.c:2891
0x609f90 lto_main()
	/home/marxin/Programming/gcc2/gcc/lto/lto.c:3356

It's logically not supported as I guess you have 2 versions of the pragma body and thus 2 variables are expected. Proper fix would be to disable the pass in case of an offloading is used.
And it's questionable whether an error should be printed for the test-case above.
Comment 6 Tom de Vries 2018-03-01 13:46:20 UTC
Looks like a duplicate of PR71536
Comment 7 Tom de Vries 2018-03-05 17:56:24 UTC
Created attachment 43571 [details]
Tentative patch

Works on nvptx for the examples from comment 0, comment 3 and comment 5.
Comment 8 Tom de Vries 2018-03-07 08:33:22 UTC
(In reply to Tom de Vries from comment #7)
> Created attachment 43571 [details]
> Tentative patch
> 
> Works on nvptx for the examples from comment 0, comment 3 and comment 5.

Tested libgomp on x86_64 with nvptx accelerator, no issues found.

Now doing bootstrap and reg-test on x86_64.
Comment 9 Tom de Vries 2018-03-07 08:33:49 UTC
(In reply to Tom de Vries from comment #6)
> Looks like a duplicate of PR71536

Actually, this is not fixed by the patch, so likely another issue.
Comment 11 Tom de Vries 2018-03-26 09:53:22 UTC
The switch conversion part of this is handled in PR85063 - "Support switch conversion in offloading functions".

Remaining to be handled here is the example from comment 5. Updating summary to reflect that.
Comment 12 Tom de Vries 2019-06-15 08:25:05 UTC
(In reply to Tom de Vries from comment #11)
> Remaining to be handled here is the example from comment 5. Updating summary
> to reflect that.

That example passes now on trunk, presumably because of r272322 ( PR90779 comment 10).
Comment 13 Tom de Vries 2019-06-15 11:51:16 UTC
Marking duplicate.

*** This bug has been marked as a duplicate of bug 90779 ***