Bug 100400 - ICE in visit_loops_in_gang_single_region
Summary: ICE in visit_loops_in_gang_single_region
Status: NEW
Alias: None
Product: gcc
Classification: Unclassified
Component: middle-end (show other bugs)
Version: 12.0
: P3 normal
Target Milestone: ---
Assignee: Not yet assigned to anyone
URL:
Keywords: openacc
Depends on:
Blocks:
 
Reported: 2021-05-03 16:53 UTC by Arseny Solokha
Modified: 2024-04-14 05:02 UTC (History)
5 users (show)

See Also:
Host:
Target:
Build:
Known to work:
Known to fail:
Last reconfirmed: 2024-04-13 00:00:00


Attachments

Note You need to log in before you can comment on or make changes to this bug.
Description Arseny Solokha 2021-05-03 16:53:30 UTC
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
Comment 1 Martin Liška 2021-05-11 12:03:01 UTC
Can't reproduce that..
Comment 2 Arseny Solokha 2021-05-11 13:22:31 UTC
% 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
Comment 3 Arseny Solokha 2021-05-11 13:26:12 UTC
% 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};
        }
    }
}
Comment 4 Martin Liška 2021-05-11 13:31:08 UTC
Sorry, my bad, I used C compiler.
Started with r11-8242-g3395dfc4da8ad1fc.
Comment 5 Thomas Schwinge 2022-01-19 20:45:18 UTC
'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.
Comment 6 GCC Commits 2022-03-04 13:25:51 UTC
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.
Comment 7 GCC Commits 2022-05-10 13:21:58 UTC
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'.
Comment 8 seurer 2022-07-15 17:47:48 UTC
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.
Comment 9 GCC Commits 2022-10-20 18:15:22 UTC
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)
Comment 10 Thomas Schwinge 2022-10-20 18:23:57 UTC
(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.