[Bug tree-optimization/100400] New: ICE in visit_loops_in_gang_single_region

[Bug tree-optimization/100400] New: ICE in visit_loops_in_gang_single_region

[asolokha at gmx dot com]

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

            Bug ID: 100400
           Summary: ICE in visit_loops_in_gang_single_region
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Keywords: openacc
          Severity: normal
          Priority: P3
         Component: tree-optimization
          Assignee: unassigned at gcc dot gnu.org
          Reporter: asolokha at gmx dot com
  Target Milestone: ---

g++-12.0.0-alpha20210502 snapshot (g:92f59e47f5a468b96b12b15233a6729904b1a1ee)
ICEs when compiling the following testcase, reduced from
libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c, w/ -O1
-fopenacc -g --param openacc-kernels=decompose:

int *p;

void
foo (void)
{
#pragma acc kernels
  {
    int c;

    p = &c;

#pragma acc loop independent
    for (c = 0; c < 1; ++c)
      ;
  }
}

% g++-12.0.0 -O1 -fopenacc -g --param openacc-kernels=decompose -c uupipzvt.c
during GIMPLE pass: omp_oacc_kernels_decompose
uupipzvt.c: In function 'void foo()':
uupipzvt.c:14:7: internal compiler error: in visit_loops_in_gang_single_region,
at omp-oacc-kernels-decompose.cc:242
   14 |       ;
      |       ^
0x8b682a visit_loops_in_gang_single_region
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:242
0x8b682a visit_loops_in_gang_single_region
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:230
0xe16a2a walk_gimple_stmt(gimple_stmt_iterator*, tree_node*
(*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**,
int*, void*), walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:578
0xe16c70 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*,
bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*),
walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:51
0xe16b39 walk_gimple_stmt(gimple_stmt_iterator*, tree_node*
(*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**,
int*, void*), walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:641
0xe16c70 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*,
bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*),
walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:51
0xe16d03 walk_gimple_seq(gimple*, tree_node* (*)(gimple_stmt_iterator*, bool*,
walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:81
0x1bf9048 make_loops_gang_single
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:290
0x1bf9048 make_region_seq
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:331
0x1bfb60c decompose_kernels_region_body
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:1336
0x1bfb60c omp_oacc_kernels_decompose_1
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:1453
0x1bfb60c omp_oacc_kernels_decompose_callback_stmt
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:1478
0xe16a2a walk_gimple_stmt(gimple_stmt_iterator*, tree_node*
(*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**,
int*, void*), walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:578
0xe16c70 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*,
bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*),
walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:51
0xe16b15 walk_gimple_stmt(gimple_stmt_iterator*, tree_node*
(*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**,
int*, void*), walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:685
0xe16c70 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*,
bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*),
walk_stmt_info*)
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/gimple-walk.c:51
0x1bf89d6 omp_oacc_kernels_decompose
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:1495
0x1bf89d6 execute
       
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210502/work/gcc-12-20210502/gcc/omp-oacc-kernels-decompose.cc:1534

[Bug tree-optimization/100400] ICE in visit_loops_in_gang_single_region

[marxin at gcc dot gnu.org]

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

Martin Liška <marxin at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
     Ever confirmed|0                           |1
             Status|UNCONFIRMED                 |WAITING
                 CC|                            |marxin at gcc dot gnu.org
   Last reconfirmed|                            |2021-05-11

--- Comment #1 from Martin Liška <marxin at gcc dot gnu.org> ---
Can't reproduce that..

[Bug tree-optimization/100400] ICE in visit_loops_in_gang_single_region

[asolokha at gmx dot com]

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

--- Comment #2 from Arseny Solokha <asolokha at gmx dot com> ---
% g++-12.0.0 -v
Using built-in specs.
COLLECT_GCC=g++-12.0.0
COLLECT_LTO_WRAPPER=/usr/libexec/gcc/x86_64-unknown-linux-gnu/12.0.0/lto-wrapper
Target: x86_64-unknown-linux-gnu
Configured with:
/var/tmp/portage/sys-devel/gcc-12.0.0_alpha20210509/work/gcc-12-20210509/configure
--host=x86_64-unknown-linux-gnu --build=x86_64-unknown-linux-gnu --prefix=/usr
--bindir=/usr/x86_64-unknown-linux-gnu/gcc-bin/12.0.0
--includedir=/usr/lib/gcc/x86_64-unknown-linux-gnu/12.0.0/include
--datadir=/usr/share/gcc-data/x86_64-unknown-linux-gnu/12.0.0
--mandir=/usr/share/gcc-data/x86_64-unknown-linux-gnu/12.0.0/man
--infodir=/usr/share/gcc-data/x86_64-unknown-linux-gnu/12.0.0/info
--with-gxx-include-dir=/usr/lib/gcc/x86_64-unknown-linux-gnu/12.0.0/include/g++-v12
--with-python-dir=/share/gcc-data/x86_64-unknown-linux-gnu/12.0.0/python
--enable-languages=c,c++ --enable-obsolete --enable-secureplt --disable-werror
--with-system-zlib --disable-nls --enable-checking=yes --disable-esp
--enable-libstdcxx-time --disable-libstdcxx-pch --enable-shared
--enable-threads=posix --enable-__cxa_atexit --enable-clocale=gnu
--disable-multilib --with-multilib-list=m64 --disable-fixed-point
--enable-targets=all --enable-libgomp --disable-libssp --disable-libada
--disable-systemtap --enable-valgrind-annotations --disable-vtable-verify
--disable-libvtv --without-zstd --enable-lto --with-isl
--disable-isl-version-check --disable-libsanitizer --enable-default-pie
--enable-default-ssp
Thread model: posix
Supported LTO compression algorithms: zlib
gcc version 12.0.0 20210509 (experimental) (GCC)

https://gcc.godbolt.org/z/3MajcK6q3

[Bug tree-optimization/100400] ICE in visit_loops_in_gang_single_region

[asolokha at gmx dot com]

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

--- Comment #3 from Arseny Solokha <asolokha at gmx dot com> ---
% cat tttt.c.008t.omp_oacc_kernels_decompose

;; Function foo (_Z3foov, funcdef_no=0, decl_uid=2487, cgraph_uid=1,
symbol_order=1)




EMERGENCY DUMP:

void foo ()
{
  # DEBUG BEGIN_STMT
  #pragma omp target oacc_kernels map(force_tofrom:p [len: 8])
    {
      int c.0;

      # DEBUG BEGIN_STMT
      try
        {
          # DEBUG BEGIN_STMT
          # DEBUG BEGIN_STMT
          p = &c;
          # DEBUG BEGIN_STMT
          #pragma acc loop independent private(c.0) private(c)
          for (c.0 = 0; c.0 < 1; c.0 = c.0 + 1)
            {
              c = c.0;
              # DEBUG BEGIN_STMT
            }
        }
      finally
        {
          c = {CLOBBER};
        }
    }
}

[Bug tree-optimization/100400] ICE in visit_loops_in_gang_single_region

[marxin at gcc dot gnu.org]

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

Martin Liška <marxin at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
             Status|WAITING                     |NEW
                 CC|                            |tschwinge at gcc dot gnu.org

--- Comment #4 from Martin Liška <marxin at gcc dot gnu.org> ---
Sorry, my bad, I used C compiler.
Started with r11-8242-g3395dfc4da8ad1fc.

[Bug middle-end/100400] ICE in visit_loops_in_gang_single_region

[tschwinge at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Last reconfirmed|2021-05-11 00:00:00         |2022-1-19
          Component|tree-optimization           |middle-end
           See Also|                            |https://gcc.gnu.org/bugzill
                   |                            |a/show_bug.cgi?id=104061,
                   |                            |https://gcc.gnu.org/bugzill
                   |                            |a/show_bug.cgi?id=103836

--- Comment #5 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
'gcc/omp-oacc-kernels-decompose.cc' is currently not at all considerate of
'GIMPLE_DEBUG' statements -- and it's not always straight forward how to handle
these (not rocket science either; but needs proper understanding and testing).

I'm assuming that all of PR100400 and PR103836 and PR104061 will be fixed at
once; but thanks for the different test cases, which do all exhibit different
characteristics.

[Bug middle-end/100400] ICE in visit_loops_in_gang_single_region

[cvs-commit at gcc dot gnu.org]

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

--- Comment #6 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Thomas Schwinge <tschwinge@gcc.gnu.org>:

https://gcc.gnu.org/g:c14ea6a72fb1ae66e3d32ac8329558497c6e4403

commit r12-7478-gc14ea6a72fb1ae66e3d32ac8329558497c6e4403
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Wed Jan 19 14:04:42 2022 +0100

    Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
[PR100400, PR103836, PR104061]

    Actually fixing it is a separate task, but it seems prudent to at least
catch
    it, and document via a few test cases.

            gcc/
            PR middle-end/100400
            PR middle-end/103836
            PR middle-end/104061
            * omp-oacc-kernels-decompose.cc (decompose_kernels_region_body):
            Catch 'GIMPLE_DEBUG'.
            gcc/testsuite/
            PR middle-end/100400
            PR middle-end/103836
            PR middle-end/104061
            * c-c++-common/goacc/kernels-decompose-pr100400-1-1.c: New.
            * c-c++-common/goacc/kernels-decompose-pr100400-1-2.c: New.
            * c-c++-common/goacc/kernels-decompose-pr100400-1-3.c: New.
            * c-c++-common/goacc/kernels-decompose-pr100400-1-4.c: New.
            * c-c++-common/goacc/kernels-decompose-pr103836-1-1.c: New.
            * c-c++-common/goacc/kernels-decompose-pr103836-1-2.c: New.
            * c-c++-common/goacc/kernels-decompose-pr103836-1-3.c: New.
            * c-c++-common/goacc/kernels-decompose-pr103836-1-4.c: New.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-1.c: New.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: New.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: New.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: New.

[Bug middle-end/100400] ICE in visit_loops_in_gang_single_region

[cvs-commit at gcc dot gnu.org]

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

--- Comment #7 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Thomas Schwinge <tschwinge@gcc.gnu.org>:

https://gcc.gnu.org/g:da6305558bab9e24943848e4fc5bd8738d7e8f9b

commit r13-262-gda6305558bab9e24943848e4fc5bd8738d7e8f9b
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Mon May 2 15:15:26 2022 +0200

    Make 'c-c++-common/goacc/kernels-decompose-pr100400-1-*.c' behave
consistently, regardless of checking level

    Fix-up for commit c14ea6a72fb1ae66e3d32ac8329558497c6e4403
    "Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
    [PR100400, PR103836, PR104061]".

    For C++ compilation of
'c-c++-common/goacc/kernels-decompose-pr100400-1-2.c',
    we first emit a 'sorry' diagnostic, and then a 'gcc_unreachable' (or
    'internal_error', see below) diagnostic, but for example, for
    '--enable-checking=release' (thus, '!CHECKING_P'), the second one may
actually
    be turned into a 'confused by earlier errors, bailing out' diagnostic. 
(See
    'gcc/diagnostic.cc:diagnostic_report_diagnostic': "When not checking, ICEs
are
    converted to fatal errors when an error has already occurred.")  Thus, make
    'c-c++-common/goacc/kernels-decompose-pr100400-1-2.c' behave consistently
via
    '-Wfatal-errors', and thus only matching the 'sorry' diagnostic.

    For example, for '--enable-checking=no' (thus, '!ENABLE_ASSERT_CHECKING'),
a
    call to 'gcc_unreachable' cannot be assumed emit an 'internal_error'-like
    diagnostic, so explicitly call 'internal_error' in
    'gcc/omp-oacc-kernels-decompose.cc:visit_loops_in_gang_single_region', in
the
    'GIMPLE_OMP_FOR' case, to avoid regressing
    'c-c++-common/goacc/kernels-decompose-pr100400-1-3.c', and
    'c-c++-common/goacc/kernels-decompose-pr100400-1-4.c'.

            PR middle-end/100400
            gcc/
            * omp-oacc-kernels-decompose.cc
            (visit_loops_in_gang_single_region) <GIMPLE_OMP_FOR>: Explicitly
            call 'internal_error'.
            gcc/testsuite/
            * c-c++-common/goacc/kernels-decompose-pr100400-1-2.c: Specify
            '-Wfatal-errors'.

[Bug middle-end/100400] ICE in visit_loops_in_gang_single_region

[seurer at gcc dot gnu.org]

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

seurer at gcc dot gnu.org changed:

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

--- Comment #8 from seurer at gcc dot gnu.org ---
Any chance we can get r13-262-gda6305558bab9e24943848e4fc5bd8738d7e8f9b
backported to gcc 12?  kernels-decompose-pr100400-1-2.c still fails there with
checks turned on.

[Bug middle-end/100400] ICE in visit_loops_in_gang_single_region

[cvs-commit at gcc dot gnu.org]

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

--- Comment #9 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The releases/gcc-12 branch has been updated by Thomas Schwinge
<tschwinge@gcc.gnu.org>:

https://gcc.gnu.org/g:d3bdc33784c857668503b30368ae50a8d2fca274

commit r12-8853-gd3bdc33784c857668503b30368ae50a8d2fca274
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Mon May 2 15:15:26 2022 +0200

    Make 'c-c++-common/goacc/kernels-decompose-pr100400-1-*.c' behave
consistently, regardless of checking level

    Fix-up for commit c14ea6a72fb1ae66e3d32ac8329558497c6e4403
    "Catch 'GIMPLE_DEBUG' misbehavior in OpenACC 'kernels' decomposition
    [PR100400, PR103836, PR104061]".

    For C++ compilation of
'c-c++-common/goacc/kernels-decompose-pr100400-1-2.c',
    we first emit a 'sorry' diagnostic, and then a 'gcc_unreachable' (or
    'internal_error', see below) diagnostic, but for example, for
    '--enable-checking=release' (thus, '!CHECKING_P'), the second one may
actually
    be turned into a 'confused by earlier errors, bailing out' diagnostic. 
(See
    'gcc/diagnostic.cc:diagnostic_report_diagnostic': "When not checking, ICEs
are
    converted to fatal errors when an error has already occurred.")  Thus, make
    'c-c++-common/goacc/kernels-decompose-pr100400-1-2.c' behave consistently
via
    '-Wfatal-errors', and thus only matching the 'sorry' diagnostic.

    For example, for '--enable-checking=no' (thus, '!ENABLE_ASSERT_CHECKING'),
a
    call to 'gcc_unreachable' cannot be assumed emit an 'internal_error'-like
    diagnostic, so explicitly call 'internal_error' in
    'gcc/omp-oacc-kernels-decompose.cc:visit_loops_in_gang_single_region', in
the
    'GIMPLE_OMP_FOR' case, to avoid regressing
    'c-c++-common/goacc/kernels-decompose-pr100400-1-3.c', and
    'c-c++-common/goacc/kernels-decompose-pr100400-1-4.c'.

            PR middle-end/100400
            gcc/
            * omp-oacc-kernels-decompose.cc
            (visit_loops_in_gang_single_region) <GIMPLE_OMP_FOR>: Explicitly
            call 'internal_error'.
            gcc/testsuite/
            * c-c++-common/goacc/kernels-decompose-pr100400-1-2.c: Specify
            '-Wfatal-errors'.

    (cherry picked from commit da6305558bab9e24943848e4fc5bd8738d7e8f9b)

[Bug middle-end/100400] ICE in visit_loops_in_gang_single_region

[tschwinge at gcc dot gnu.org]

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

--- Comment #10 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
(In reply to seurer from comment #8)
> Any chance we can get r13-262-gda6305558bab9e24943848e4fc5bd8738d7e8f9b
> backported to gcc 12?  kernels-decompose-pr100400-1-2.c still fails there

Now done, and sorry for the delay.

---

This PR remains open for the underlying issue has not yet been addressed.