[Bug target/104968] New: [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[Bug target/104968] New: [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[burnus at gcc dot gnu.org]

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

[Bug target/104968] [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104968

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
Can't reproduce.

It this not fixed by:
...
commit 7862f6ccd85a001e4d70abb00bb95d8c7846ba80
Author: Tom de Vries <tdevries@suse.de>
Date:   Wed Feb 23 09:33:33 2022 +0100

    [nvptx] Fix dummy location in gen_comment
...
?

[Bug target/104968] [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104968

--- Comment #2 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #1)
> Can't reproduce.
> 
> It this not fixed by:
> ...
> commit 7862f6ccd85a001e4d70abb00bb95d8c7846ba80
> Author: Tom de Vries <tdevries@suse.de>
> Date:   Wed Feb 23 09:33:33 2022 +0100
> 
>     [nvptx] Fix dummy location in gen_comment
> ...
> ?

Hmm, wait, of course I have a patch in my stack that's pending for upstream.
Let me undo that one and retry.

[Bug target/104968] [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104968

--- Comment #3 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #2)
> (In reply to Tom de Vries from comment #1)
> > Can't reproduce.
> > 
> > It this not fixed by:
> > ...
> > commit 7862f6ccd85a001e4d70abb00bb95d8c7846ba80
> > Author: Tom de Vries <tdevries@suse.de>
> > Date:   Wed Feb 23 09:33:33 2022 +0100
> > 
> >     [nvptx] Fix dummy location in gen_comment
> > ...
> > ?
> 
> Hmm, wait, of course I have a patch in my stack that's pending for upstream.
> Let me undo that one and retry.

Ack, reproduced.

[Bug target/104968] [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104968

--- Comment #4 from Tom de Vries <vries at gcc dot gnu.org> ---
This ( https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591912.html )
proposed patch fixes this ICE, pinged again.

[Bug target/104968] [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104968

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |jakub at gcc dot gnu.org

--- Comment #5 from Tom de Vries <vries at gcc dot gnu.org> ---
This patch fixes the ICE at openmp level:
...
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 139a0de6100..19af384c634 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -13361,6 +13361,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
       g = gimple_build_bind (NULL_TREE, gfor, NULL_TREE);
       g = gimple_build_omp_task (g, task_clauses, NULL_TREE, NULL_TREE,
                                 NULL_TREE, NULL_TREE, NULL_TREE);
+      gimple_set_location (g, EXPR_LOCATION (*expr_p));
       gimple_omp_task_set_taskloop_p (g, true);
       g = gimple_build_bind (NULL_TREE, g, NULL_TREE);
       gomp_for *gforo
...

[Bug target/104968] [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104968

--- Comment #6 from Tom de Vries <vries at gcc dot gnu.org> ---
(In reply to Tom de Vries from comment #5)
> This patch fixes the ICE at openmp level:
> ...
> diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
> index 139a0de6100..19af384c634 100644
> --- a/gcc/gimplify.cc
> +++ b/gcc/gimplify.cc
> @@ -13361,6 +13361,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
>        g = gimple_build_bind (NULL_TREE, gfor, NULL_TREE);
>        g = gimple_build_omp_task (g, task_clauses, NULL_TREE, NULL_TREE,
>                                  NULL_TREE, NULL_TREE, NULL_TREE);
> +      gimple_set_location (g, EXPR_LOCATION (*expr_p));
>        gimple_omp_task_set_taskloop_p (g, true);
>        g = gimple_build_bind (NULL_TREE, g, NULL_TREE);
>        gomp_for *gforo
> ...

Submitted a more complete patch here (
https://gcc.gnu.org/pipermail/gcc-patches/2022-March/591954.html ).

[Bug target/104968] [nvptx][OpenMP] SIGSEGV / ICE in final_scan_insn_1

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104968

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |FIXED
   Target Milestone|---                         |12.0
             Status|UNCONFIRMED                 |RESOLVED

--- Comment #7 from Tom de Vries <vries at gcc dot gnu.org> ---
Fixed by
https://gcc.gnu.org/git/?p=gcc.git;a=commit;h=356e2720e9030927579024c2f060d665a0b9080f
.