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.
Can you please help me how to run the single test on my machine w/o nvptx target?
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
(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.
Good, now I can reproduce that with intelmic offloading compiler.
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.
Looks like a duplicate of PR71536
Created attachment 43571 [details] Tentative patch Works on nvptx for the examples from comment 0, comment 3 and comment 5.
(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.
(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.
https://gcc.gnu.org/ml/gcc-patches/2018-03/msg00320.html
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.
(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).
Marking duplicate. *** This bug has been marked as a duplicate of bug 90779 ***