https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104968
Bug ID: 104968
Summary: [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1
Product: gcc
Version: 12.0
Status: UNCONFIRMED
Keywords: ice-on-valid-code, openmp
Severity: normal
Priority: P3
Component: target
Assignee: unassigned at gcc dot gnu.org
Reporter: burnus at gcc dot gnu.org
CC: vries at gcc dot gnu.org
Target Milestone: ---
Target: nvptx-none
* Testcase is
https://github.com/clang-ykt/omptests/blob/master/t-taskloopsimd/test.c
* Reduced testcase: see below.
In either case:
* When compiled with -fopenmp -O1, it fails as follows. With -O0, it works.
This seems to be a REGRESSION:
* it PASSES on OG11 (devel/omp/gcc-11) branch
(which should have no local patches in this area)
* it fails on mainline/GCC 12
With our approx. bi-nightly tester, I see PASS, an intermittend fail on
2022-01-11,
then PASS until 2022-02-20 and FAIL since 2022-02-22.
In that time range, I see the following commit (but it might also be
caused/exposed by another commit):
commit g:5ed77fb3ed1ee0289a0ec9499ef52b99b39421f1
CommitDate: Tue Feb 22 15:48:03 2022 +0100
[libgomp, nvptx] Fix hang in gomp_team_barrier_wait_end
commit g:c2b23aaaf4457278403c01cd145cd3936683384e
CommitDate: Tue Feb 22 14:51:59 2022 +0100
[nvptx] Add -mptx-comment
commit g:02aedc6f269b5e3c1f354edcf5b84d27b0a15946
CommitDate: Mon Feb 21 16:49:37 2022 +0100
[nvptx] Initialize ptx regs
commit 69cb3f2abb911acebfc7ffede2ee7151a3e14a59
CommitDate: Sat Feb 19 20:05:56 2022 +0100
[nvptx] Use _ as destination operand of atom.exch
Program received signal SIGSEGV, Segmentation fault.
final_scan_insn_1 (insn=0x7ffff743ecc0, file=<optimized out>,
seen=0x7fffffffddbc, nopeepholes=<optimized out>, optimize_p=<optimized out>)
at gcc-mainline/gcc/final.cc:2648
2648 if (*loc.file && loc.line)
(gdb) bt
#0 final_scan_insn_1 (insn=0x7ffff743ecc0, file=<optimized out>,
seen=0x7fffffffddbc, nopeepholes=<optimized out>, optimize_p=<optimized out>)
at gcc-mainline/gcc/final.cc:2648
#1 0x000000000087eeac in final_scan_insn (insn=<optimized out>,
file=<optimized out>, optimize_p=<optimized out>, nopeepholes=<optimized out>,
seen=<optimized out>)
at gcc-mainline/gcc/final.cc:2940
#2 0x000000000087efc6 in final_1 (first=0x7ffff742b460, file=0x21c59e0,
seen=1, optimize_p=1) at gcc-mainline/gcc/final.cc:1997
#3 0x000000000087ffe5 in rest_of_handle_final () at
gcc-mainline/gcc/final.cc:4285
#4 (anonymous namespace)::pass_final::execute (this=<optimized out>) at
gcc-mainline/gcc/final.cc:4363
#5 0x0000000000b97a2b in execute_one_pass (pass=pass@entry=0x21eeca0) at
gcc-mainline/gcc/passes.cc:2637
#6 0x0000000000b98398 in execute_pass_list_1 (pass=0x21eeca0) at
gcc/passes.cc:2737
#7 0x0000000000b983aa in execute_pass_list_1 (pass=0x21ee7c0) at
gcc/passes.cc:2738
#8 0x0000000000b983aa in execute_pass_list_1 (pass=0x21eccc0) at
gcc/passes.cc:2738
#9 0x0000000000b983f5 in execute_pass_list (fn=<optimized out>,
pass=<optimized out>) at gcc-mainline/gcc/passes.cc:2748
#10 0x000000000076423d in cgraph_node::expand (this=0x7ffff73f7550) at
gcc-mainline/gcc/cgraphunit.cc:1834
#11 0x000000000076599d in expand_all_functions () at
gcc-mainline/gcc/cgraphunit.cc:1998
Doing in the debugger 'up', 'up' (as isns are optimized out) and then the
following:
(gdb) p debug_rtx(insn)
(insn 113 115 112 14 (asm_input ("// Start: Added by -minit-regs=3:")) -1
(nil))
==== Reduced testcase ====
int main()
{
double a[10], a_h[10];
int myId = -1;
#pragma omp target map(tofrom:a)
#pragma omp taskloop simd shared(a) lastprivate(myId)
for(int i = 0 ; i < 10; i++) if (a[i] != a_h[i]) { }
}
