[Bug middle-end/104757] New: [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[Bug middle-end/104757] New: [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[burnus at gcc dot gnu.org]

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

            Bug ID: 104757
           Summary: [12 Regression] ICE (segfault) GIMPLE pass: walloca -
                    in gimple_range_global
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Keywords: ice-on-valid-code
          Severity: normal
          Priority: P3
         Component: middle-end
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
  Target Milestone: ---

This requires that
   ENABLE_OFFLOADING
is true (affects gcc/*.cc; e.g. by s/ENABLE_OFFLOADING/true/ inside gcc/*.cc).

Plus the following environment variable has to be set
OFFLOAD_TARGET_NAMES=nvptx-none
(This value, using something else [like amdgcn-amdhsa] does not ICE.)


Afterwards, the fail is:

  export OFFLOAD_TARGET_NAMES=nvptx-none
  f951 -quiet -O -fopenmp testsuite/gfortran.dg/gomp/clauses-1.f90

during GIMPLE pass: walloca
testsuite/gfortran.dg/gomp/clauses-1.f90:61:14:

   61 |     ll = ll +1
      |              ^
internal compiler error: Segmentation fault
0x100e603 crash_signal
        ../../repos/gcc/gcc/toplev.cc:322
0x13246be gimple_range_global(tree_node*)
        ../../repos/gcc/gcc/value-query.cc:424
0x1d2d910 ranger_cache::get_global_range(irange&, tree_node*) const
        ../../repos/gcc/gcc/gimple-range-cache.cc:923
0x1d2d910 ranger_cache::get_global_range(irange&, tree_node*) const

[Bug middle-end/104757] [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[tschwinge at gcc dot gnu.org]

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

Thomas Schwinge <tschwinge at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|UNCONFIRMED                 |NEW
   Last reconfirmed|                            |2022-03-02
                 CC|                            |tschwinge at gcc dot gnu.org
     Ever confirmed|0                           |1

--- Comment #1 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Confirmed ever since commit 48c6cac9caea1dc7c5f50ad3a736f6693e74a11b
"Fortran/openmp: Fix '!$omp end'", but only in an offloading-enabled build but
not without offloading.  (Specifically nvptx offloading, as you've now
determined.)

[Bug middle-end/104757] [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[burnus at gcc dot gnu.org]

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

--- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
The reason for the nvptx dependency is the following.

omp_max_simt_vf (void)
{
  if (!optimize)
    return 0;
  if (ENABLE_OFFLOADING)
    for (const char *c = getenv ("OFFLOAD_TARGET_NAMES"); c;)
      {
        if (startswith (c, "nvptx"))
          return 32;
...
  return 0;
}

 * * *

Reduced testcase:

module m
  implicit none (external, type)
  integer :: ll
  !$omp declare target (ll)
contains
subroutine foo (i1)
  !$omp declare target (foo)
  logical :: i1
  integer :: i
  !$omp distribute simd if(i1)
  do i = 1, 64
    ll = ll +1
  end do
end
end module

[Bug middle-end/104757] [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[burnus at gcc dot gnu.org]

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

--- Comment #3 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Thomas Schwinge from comment #1)
> Confirmed ever since commit 48c6cac9caea1dc7c5f50ad3a736f6693e74a11b
which was committed on 'Nov 12 17:58:21 2021'. — However, that's a FE-only
patch which only affects parsing.

 * * *

The ICE/segfault is in  ranger_cache::get_global_range (irange &r, tree name)
const   for
923       r = gimple_range_global (name);

(gdb) p debug_tree(name)
 <ssa_name 0x7ffff6fd1798
    type <boolean_type 0x7ffff6fed150 logical(kind=4) public unsigned SI
    ...

which is a 4-byte Boolean value.

[Bug middle-end/104757] [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[rguenth at gcc dot gnu.org]

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

--- Comment #4 from Richard Biener <rguenth at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #3)
> (In reply to Thomas Schwinge from comment #1)
> > Confirmed ever since commit 48c6cac9caea1dc7c5f50ad3a736f6693e74a11b
> which was committed on 'Nov 12 17:58:21 2021'. — However, that's a FE-only
> patch which only affects parsing.
> 
>  * * *
> 
> The ICE/segfault is in  ranger_cache::get_global_range (irange &r, tree
> name) const   for
> 923       r = gimple_range_global (name);
> 
> (gdb) p debug_tree(name)
>  <ssa_name 0x7ffff6fd1798
>     type <boolean_type 0x7ffff6fed150 logical(kind=4) public unsigned SI
>     ...
> 
> which is a 4-byte Boolean value.

you don't quote more but I guess this name is released and thus has a NULL
SSA_NAME_DEF_STMT so the is_a check crashes.

  if (SSA_NAME_IS_DEFAULT_DEF (name) || (cfun && cfun->after_inlining)
      || is_a<gphi *> (SSA_NAME_DEF_STMT (name)))

[Bug middle-end/104757] [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[rguenth at gcc dot gnu.org]

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

Richard Biener <rguenth at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|---                         |12.0
                 CC|                            |rguenth at gcc dot gnu.org

--- Comment #5 from Richard Biener <rguenth at gcc dot gnu.org> ---
The backtrace is also too short to tell anything meaningful.

[Bug middle-end/104757] [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[burnus at gcc dot gnu.org]

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

--- Comment #6 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Tobias Burnus from comment #3)
> (In reply to Thomas Schwinge from comment #1)
> > Confirmed ever since commit 48c6cac9caea1dc7c5f50ad3a736f6693e74a11b
... which added the testcase (by porting it from the C testcase).

Testing with the reduced testcase shows:

* FAILS since r12-5504-gd986ff50b4aad62c45d7ac62915e072643ddfca1
  "Split return functionality of get_non_stale_global_range." (Nov 24, 2021)
  with the ICE mentioned before:
    during GIMPLE pass: walloca  in gimple_range_global

Before that version, it also gave an ICE - but at a different place:

* FAILS since r12-2891-gd0befed793b94f3f407be44e6f69f81a02f5f073
  "openmp: Add support for OpenMP 5.1 masked construct" (Aug 12, 2021)
  with the ICE
     during GIMPLE pass: ssa
     input.f90:15:3: internal compiler error: Segmentation fault
     15 | end
        |   ^
     0xf6e773 crash_signal
            ../../repos/gcc/gcc/toplev.c:328
     0x11b9652 verify_ssa(bool, bool)
         ../../repos/gcc/gcc/tree-ssa.c:1076
     0xe7ecf5 execute_function_todo
        ../../repos/gcc/gcc/passes.c:2049

* Before that commit and with the current GCC 11 branch, it works.

 * * *

Current-mainline backtrace:

during GIMPLE pass: walloca
gcc/foo.f90:13:14:

   13 |     ll = ll +1
      |              ^
internal compiler error: Segmentation fault
0x100e6d3 crash_signal
        ../../repos/gcc/gcc/toplev.cc:322
0x132478e gimple_range_global(tree_node*)
        ../../repos/gcc/gcc/value-query.cc:424
0x1d2d9e0 ranger_cache::get_global_range(irange&, tree_node*) const
        ../../repos/gcc/gcc/gimple-range-cache.cc:923
0x1d2d9e0 ranger_cache::get_global_range(irange&, tree_node*) const
        ../../repos/gcc/gcc/gimple-range-cache.cc:919
0x1d2975a gimple_ranger::export_global_ranges()
        ../../repos/gcc/gcc/gimple-range.cc:474
0x1d6bdbc pass_walloca::execute(function*)
        ../../repos/gcc/gcc/gimple-ssa-warn-alloca.cc:381

Or in the debugger:

0x000000000132478e in gimple_range_global (name=name@entry=0x7ffff6fd1798) at
../../repos/gcc/gcc/gimple.h:1243
1243      return gs->code == GIMPLE_PHI;

(gdb) bt
#0  0x000000000132478e in gimple_range_global (name=name@entry=0x7ffff6fd1798)
at ../../repos/gcc/gcc/gimple.h:1243
#1  0x0000000001d2d9e1 in ranger_cache::get_global_range (this=0x28,
name=0x7ffff6fd1798, r=...) at ../../repos/gcc/gcc/gimple-range-cache.cc:923
#2  ranger_cache::get_global_range (this=this@entry=0x2f09288, r=...,
name=name@entry=0x7ffff6fd1798) at
../../repos/gcc/gcc/gimple-range-cache.cc:919
#3  0x0000000001d2975b in gimple_ranger::export_global_ranges
(this=this@entry=0x2f09260) at ../../repos/gcc/gcc/gimple-range-fold.h:98
#4  0x0000000001d6bdbd in pass_walloca::execute (this=0x2ee9050,
fun=0x7ffff71be000) at ../../repos/gcc/gcc/gimple-ssa-warn-alloca.cc:381
#5  0x0000000000f1c1bc in execute_one_pass (pass=0x2ee9050) at
../../repos/gcc/gcc/passes.cc:2637
#6  0x0000000000f1c980 in execute_pass_list_1 (pass=0x2ee9050) at
../../repos/gcc/gcc/passes.cc:2737
#7  0x0000000000f1c9bd in execute_pass_list (fn=0x7ffff71be000, pass=<optimized
out>) at ../../repos/gcc/gcc/passes.cc:2748
#8  0x0000000000b4d379 in cgraph_node::analyze (this=0x7ffff71c1000) at
../../repos/gcc/gcc/context.h:48
#9  0x0000000000b50592 in analyze_functions (first_time=<optimized out>) at
../../repos/gcc/gcc/cgraphunit.cc:1240
#10 0x0000000000b51082 in symbol_table::finalize_compilation_unit
(this=0x7ffff6fcd000) at ../../repos/gcc/gcc/cgraphunit.cc:2500
#11 0x000000000100e9f4 in compile_file () at ../../repos/gcc/gcc/toplev.cc:479
#12 0x0000000000919e1c in do_compile (no_backend=false) at
../../repos/gcc/gcc/toplev.cc:2168

[Bug middle-end/104757] [12 Regression] ICE (segfault) GIMPLE pass: walloca - in gimple_range_global

[burnus at gcc dot gnu.org]

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

--- Comment #7 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Richard Biener from comment #4)
> (In reply to Tobias Burnus from comment #3)
> > 923       r = gimple_range_global (name);
> > 
> > (gdb) p debug_tree(name)
> >  <ssa_name 0x7ffff6fd1798
> >     type <boolean_type 0x7ffff6fed150 logical(kind=4) public unsigned SI
> >     ...
> you don't quote more but I guess this name is released and thus has a NULL
> SSA_NAME_DEF_STMT so the is_a check crashes.
> 
>   if (SSA_NAME_IS_DEFAULT_DEF (name) || (cfun && cfun->after_inlining)
>       || is_a<gphi *> (SSA_NAME_DEF_STMT (name)))

Indeed:

(gdb) p name->base.default_def_flag
$2 = 0
(gdb) p cfun
$3 = (function *) 0x7ffff71be000
(gdb) p cfun->after_inlining 
$4 = 0
(gdb) p name->ssa_name.def_stmt
$5 = (gimple *) 0x0

[Bug middle-end/104757] [12 Regression][OpenMP] ICE in GIMPLE pass: walloca - in gimple_range_global / segfault as SSA_NAME_DEF_STMT is NULL for 'if' clause arg

[jakub at gcc dot gnu.org]

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

Jakub Jelinek <jakub at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Assignee|unassigned at gcc dot gnu.org      |jakub at gcc dot gnu.org
             Status|NEW                         |ASSIGNED

--- Comment #8 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
Created attachment 52551
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=52551&action=edit
gcc12-pr104757.patch

Untested fix.  The other option would be to make sure the SIMD -> SIMD + SIMT
duplication doesn't duplicate SSA_NAMEs that are defined outside of the
construct.

[Bug middle-end/104757] [12 Regression][OpenMP] ICE in GIMPLE pass: walloca - in gimple_range_global / segfault as SSA_NAME_DEF_STMT is NULL for 'if' clause arg

[jakub at gcc dot gnu.org]

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

--- Comment #9 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
It isn't sufficient though, consider:
#pragma omp declare target
void
foo (int x, int y, int *z)
{
  int j = 0;
  #pragma omp simd linear(j:x + y)
  for (int i = 0; i < 64; i++)
    j += x + y;
}
#pragma omp end declare target

-O2 -fopenmp with offloading to nvptx.

This is quite hard to deal with in scan_omp_simd, that calls
copy_gimple_seq_and_replace_locals and we do want to remap some stuff even in
the outermost construct's clauses, but apparently not all.
And copy_gimple_seq_and_replace_locals calls replace_locals_op which for
VAR_DECLs etc. checks if we have seen them in bind vars etc., but for SSA_NAMEs
always remaps them.

[Bug middle-end/104757] [12 Regression][OpenMP] ICE in GIMPLE pass: walloca - in gimple_range_global / segfault as SSA_NAME_DEF_STMT is NULL for 'if' clause arg

[jakub at gcc dot gnu.org]

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

Jakub Jelinek <jakub at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
  Attachment #52551|0                           |1
        is obsolete|                            |

--- Comment #10 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
Created attachment 52552
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=52552&action=edit
gcc12-pr104757.patch

Updated fix.

[Bug middle-end/104757] [12 Regression][OpenMP] ICE in GIMPLE pass: walloca - in gimple_range_global / segfault as SSA_NAME_DEF_STMT is NULL for 'if' clause arg

[cvs-commit at gcc dot gnu.org]

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

--- Comment #11 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Jakub Jelinek <jakub@gcc.gnu.org>:

https://gcc.gnu.org/g:431414b5d934866af3f6415a56c35bb57b928fef

commit r12-7459-g431414b5d934866af3f6415a56c35bb57b928fef
Author: Jakub Jelinek <jakub@redhat.com>
Date:   Thu Mar 3 09:13:32 2022 +0100

    openmp: Disable SSA form during gimplification on OMP_SIMD clauses and body
[PR104757]

    When offloading to nvptx is enabled, scan_omp_simd duplicates the simd
    region including its clauses and body using inliner's
    copy_gimple_seq_and_replace_locals.  That works nicely for decls, remaps
    only those that are seen in the nested bind expr vars (i.e. local
variables)
    and doesn't remap other vars.  But for SSA_NAMEs it remaps them always,
doesn't
    know if their def stmt is outside of the simd (then it better shouldn't be
remapped)
    or inside of it (then it should) and without cfg/dominators that is pretty
hard
    to figure out (well, we could walk the region twice, once note SSA_NAMEs
defined
    by each stmt seen there and once do the remapping of only those visited
SSA_NAMEs).

    This patch uses a simpler way, disables temporarily into_ssa for the
clauses and
    body of each simd region; we already disable into_ssa e.g. in
parallel/target/task
    etc. regions through push_gimplify_context () but for simd we don't push
    any gimplification context and appart from into_ssa I think we don't need
it.

    2022-03-03  Jakub Jelinek  <jakub@redhat.com>

            PR middle-end/104757
            * gimplify.cc (gimplify_omp_loop): Call gimplify_expr rather than
            gimplify_omp_for.
            (gimplify_expr) <case OMP_SIMD>: Temporarily disable
            gimplify_ctxp->into_ssa around call to gimplify_omp_for.

            * gfortran.dg/gomp/pr104757.f90: New test.
            * gcc.dg/gomp/pr104757.c: New test.

[Bug middle-end/104757] [12 Regression][OpenMP] ICE in GIMPLE pass: walloca - in gimple_range_global / segfault as SSA_NAME_DEF_STMT is NULL for 'if' clause arg

[jakub at gcc dot gnu.org]

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

--- Comment #12 from Jakub Jelinek <jakub at gcc dot gnu.org> ---
I admit I haven't built an offloading compiler, just hand edited auto-host.h to
ENABLE_OFFLOADING 1 and used the env var to debug those 2 testcases and then
tested on non-offloading compiler.
So, not closing this just yet until confirmed it works for offloading compiler
too.

[Bug middle-end/104757] [12 Regression][OpenMP] ICE in GIMPLE pass: walloca - in gimple_range_global / segfault as SSA_NAME_DEF_STMT is NULL for 'if' clause arg

[tschwinge at gcc dot gnu.org]

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

Thomas Schwinge <tschwinge at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|ASSIGNED                    |RESOLVED
         Resolution|---                         |FIXED

--- Comment #13 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
(In reply to Jakub Jelinek from comment #12)
> not closing this just yet until confirmed it works for offloading compiler too.

ACK, thanks.  New test cases PASS, and additionally the expected:

    [-FAIL: gfortran.dg/gomp/clauses-1.f90   -O  (internal compiler error:
Segmentation fault)-]
    [-FAIL:-]{+PASS:+} gfortran.dg/gomp/clauses-1.f90   -O  (test for excess
errors)