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
Can't reproduce that..
% 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
% 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}; } } }
Sorry, my bad, I used C compiler. Started with r11-8242-g3395dfc4da8ad1fc.
'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.
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.
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'.
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.
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)
(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.