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]) { } }