[Bug tree-optimization/100280] New: ICE in lower_omp_target, at omp-low.c:12287

[Bug tree-optimization/100280] New: ICE in lower_omp_target, at omp-low.c:12287

[asolokha at gmx dot com]

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

            Bug ID: 100280
           Summary: ICE in lower_omp_target, at omp-low.c:12287
           Product: gcc
           Version: 11.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: ---

gcc-11.0.1-alpha20210426 snapshot (g:d3212299e2cfc3c16dd23bab26ec6c49024105f8)
ICEs when compiling the following testcase, reduced from
libgomp/testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c, w/ -fopenacc
--param openacc-kernels=decompose:

void
foo (void)
{
#pragma acc kernels
  for (int i;;)
    ;
}

% gcc-11.0.1 -fopenacc --param openacc-kernels=decompose -c bz3qiqjv.c
during GIMPLE pass: omplower
bz3qiqjv.c: In function 'foo':
bz3qiqjv.c:4:9: internal compiler error: in lower_omp_target, at
omp-low.c:12287
    4 | #pragma acc kernels
      |         ^~~
0x6a5d35 lower_omp_target
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:12287
0x6a5d35 lower_omp_1
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13354
0x6a5d35 lower_omp
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13492
0xcca59e lower_omp_1
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13289
0xcca59e lower_omp
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13492
0xcc9d4e lower_omp_1
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13281
0xcc9d4e lower_omp
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13492
0xccb283 lower_omp_target
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:12138
0xccb283 lower_omp_1
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13354
0xccb283 lower_omp
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13492
0xcca59e lower_omp_1
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13289
0xcca59e lower_omp
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13492
0xcd3205 execute_lower_omp
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13531
0xcd3205 execute
       
/var/tmp/portage/sys-devel/gcc-11.0.1_alpha20210426/work/gcc-11-20210426/gcc/omp-low.c:13579

[Bug tree-optimization/100280] ICE in lower_omp_target, at omp-low.c:12287

[cvs-commit at gcc dot gnu.org]

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

--- Comment #1 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:9b32c1669aad5459dd053424f9967011348add83

commit r12-6542-g9b32c1669aad5459dd053424f9967011348add83
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Thu Dec 16 22:02:37 2021 +0100

    OpenACC 'kernels' decomposition: Mark variables used in synthesized data
clauses as addressable [PR100280]

    ... as otherwise 'gcc/omp-low.c:lower_omp_target' has to create a
temporary:

        13073                       else if (is_gimple_reg (var))
        13074                         {
        13075                           gcc_assert (offloaded);
        13076                           tree avar = create_tmp_var (TREE_TYPE
(var));
        13077                           mark_addressable (avar);

    ..., which (a) is only implemented for actualy *offloaded* regions (but not
    data regions), and (b) the subsequently synthesized code for writing to and
    later reading back from the temporary fundamentally conflicts with OpenACC
    'async' (as used by OpenACC 'kernels' decomposition).  That's all not
trivial
    to make work, so let's just avoid this case.

            gcc/
            PR middle-end/100280
            * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region):
            Mark variables used in synthesized data clauses as addressable.
            gcc/testsuite/
            PR middle-end/100280
            * c-c++-common/goacc/kernels-decompose-pr100280-1.c: New.
            * c-c++-common/goacc/classify-kernels-parloops.c: Likewise.
            * c-c++-common/goacc/classify-kernels-unparallelized-parloops.c:
            Likewise.
            * c-c++-common/goacc/classify-kernels-unparallelized.c: Test
            '--param openacc-kernels=decompose'.
            * c-c++-common/goacc/classify-kernels.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-2.c: Update.
            * c-c++-common/goacc/kernels-decompose-ice-1.c: Remove.
            * c-c++-common/goacc/kernels-decompose-ice-2.c: Likewise.
            * gfortran.dg/goacc/classify-kernels-parloops.f95: New.
            * gfortran.dg/goacc/classify-kernels-unparallelized-parloops.f95:
            Likewise.
            * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Test
            '--param openacc-kernels=decompose'.
            * gfortran.dg/goacc/classify-kernels.f95: Likewise.
            libgomp/
            PR middle-end/100280
            *
testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c:
            Update.
            * testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: Likewise.
            * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
            Likewise.

    Suggested-by: Julian Brown <julian@codesourcery.com>

[Bug tree-optimization/100280] ICE in lower_omp_target, at omp-low.c:12287

[tschwinge at gcc dot gnu.org]

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

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

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |jules at gcc dot gnu.org,
                   |                            |tschwinge at gcc dot gnu.org
         Resolution|---                         |FIXED
           See Also|https://gcc.gnu.org/bugzill |
                   |a/show_bug.cgi?id=100400    |
           Assignee|unassigned at gcc dot gnu.org      |tschwinge at gcc dot gnu.org
             Status|UNCONFIRMED                 |RESOLVED

--- Comment #2 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Thanks for the report, now fixed in master branch.  Not planning on backporting
OpenACC 'kernels' decomposition changes to release branches -- unless that'd be
useful for you?

[Bug tree-optimization/100280] ICE in lower_omp_target, at omp-low.c:12287

[asolokha at gmx dot com]

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

--- Comment #3 from Arseny Solokha <asolokha at gmx dot com> ---
(In reply to Thomas Schwinge from comment #2)
> Thanks for the report, now fixed in master branch.  Not planning on
> backporting OpenACC 'kernels' decomposition changes to release branches --
> unless that'd be useful for you?

Thanks. Not being OpenACC user, I'm fine with that.

[Bug tree-optimization/100280] ICE in lower_omp_target, at omp-low.c:12287

[cvs-commit at gcc dot gnu.org]

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

--- Comment #4 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:e5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc

commit r12-7479-ge5ae22c56152b1a1f4b4e1d7ae04431a9e4710cc
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Tue Feb 15 16:54:30 2022 +0100

    Add diagnostic: "note: OpenACC 'kernels' decomposition: variable '[...]'
declared in block made addressable" [PR100280]

    Follow-up to commit 9b32c1669aad5459dd053424f9967011348add83
    "OpenACC 'kernels' decomposition: Mark variables used in
    synthesized data clauses as addressable [PR100280]".

            PR middle-end/100280
            gcc/
            * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region):
            Add diagnostic: "note: OpenACC 'kernels' decomposition: variable
            '[...]' declared in block made addressable".
            gcc/testsuite/
            * c-c++-common/goacc/classify-kernels-unparallelized.c: Add
            '--param=openacc-privatization=noisy'.
            * c-c++-common/goacc/classify-kernels.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-2.c: Adjust.
            * c-c++-common/goacc/kernels-decompose-pr100280-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise.
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: Adjust.
            * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
            Likewise.

[Bug tree-optimization/100280] ICE in lower_omp_target, at omp-low.c:12287

[cvs-commit at gcc dot gnu.org]

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

--- Comment #5 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:de6e81ea961219d0726db67776d11ce75a4cae1b

commit r12-7480-gde6e81ea961219d0726db67776d11ce75a4cae1b
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Tue Feb 15 23:03:49 2022 +0100

    OpenACC 'kernels' decomposition: Move 'TREE_ADDRESSABLE' setting into OMP
lowering [PR100280]

    ... in preparation for later changes.  No functional change.

    Follow-up to commit 9b32c1669aad5459dd053424f9967011348add83
    "OpenACC 'kernels' decomposition: Mark variables used in
    synthesized data clauses as addressable [PR100280]".

            PR middle-end/100280
            gcc/
            * tree.h (OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE): New.
            * tree-core.h: Document it.
            * omp-low.cc (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Handle
            'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE'.
            * omp-oacc-kernels-decompose.cc (maybe_build_inner_data_region):
            Set 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' instead of
            'TREE_ADDRESSABLE'.
            gcc/testsuite/
            * c-c++-common/goacc/classify-kernels-unparallelized.c: Adjust.
            * c-c++-common/goacc/classify-kernels.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-2.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr100280-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise.
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/f-asyncwait-1.c: Adjust.
            * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
            Likewise.

[Bug tree-optimization/100280] ICE in lower_omp_target, at omp-low.c:12287

[cvs-commit at gcc dot gnu.org]

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

--- 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:8935589b496f755e08cadf26d8ceddf0dd6e0968

commit r12-7481-g8935589b496f755e08cadf26d8ceddf0dd6e0968
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Tue Feb 15 23:31:34 2022 +0100

    OMP lowering: Regimplify 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs
[PR100280, PR104132, PR104133]

    ... by generalizing the existing 'gcc/omp-low.cc:task_shared_vars'.

    Fix-up for commit 9b32c1669aad5459dd053424f9967011348add83
    "OpenACC 'kernels' decomposition: Mark variables used in
    synthesized data clauses as addressable [PR100280]".

            PR middle-end/100280
            PR middle-end/104132
            PR middle-end/104133
            gcc/
            * omp-low.cc (task_shared_vars): Rename to
            'make_addressable_vars'.  Adjust all users.
            (scan_sharing_clauses) <OMP_CLAUSE_MAP> Use it for
            'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE' DECLs, too.
            gcc/testsuite/
            * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Adjust.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise.
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
            Extend.

[Bug tree-optimization/100280] ICE in lower_omp_target, at omp-low.c:12287

[cvs-commit at gcc dot gnu.org]

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

--- 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:337ed336d7dd83526891bdb436f0bfe9e351f69d

commit r12-7624-g337ed336d7dd83526891bdb436f0bfe9e351f69d
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Thu Feb 17 14:18:57 2022 +0100

    OpenACC 'kernels' decomposition: Mark variables used in 'present' clauses
as addressable [PR100280, PR104086]

    ... like in recent commit 9b32c1669aad5459dd053424f9967011348add83
    "OpenACC 'kernels' decomposition: Mark variables used in synthesized
    data clauses as addressable [PR100280]".  Otherwise, we may run into
    'gcc/omp-low.cc:lower_omp_target':

        13125                       else if (is_gimple_reg (var))
        13126                         {
        13127                           gcc_assert (offloaded);

            PR middle-end/100280
            PR middle-end/104086
            gcc/
            * omp-oacc-kernels-decompose.cc (omp_oacc_kernels_decompose_1):
            Mark variables used in 'present' clauses as addressable.
            * omp-low.cc (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Gracefully
            handle duplicate 'OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE'.
            gcc/testsuite/
            * c-c++-common/goacc/kernels-decompose-pr104086-1.c: Adjust,
            extend.
            libgomp/
            *
testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c:
            Merge this...
            *
testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c:
            ..., and this...
            * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: ... into
            this, and adjust.
            * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
            Extend.

[Bug tree-optimization/100280] ICE in lower_omp_target, at omp-low.c:12287

[cvs-commit at gcc dot gnu.org]

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

--- Comment #8 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:a07b8f4fb756484893b5612cbe9410970dc76db9

commit r12-7627-ga07b8f4fb756484893b5612cbe9410970dc76db9
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Fri Mar 11 22:31:51 2022 +0100

    OpenACC 'kernels' decomposition: resolve wrong-code cases unless manually
making certain variables addressable [PR100280, PR104892]

    Currently in OpenACC 'kernels' decomposition, there is special handling of
    'GOMP_MAP_FORCE_TOFROM', documented to be done to avoid "internal compiler
    errors in later passes".  For performance reasons, the current repetitive
    to/from device copying for every region is not ideal, compared to using
    'present' clauses, as done for almost all other 'GOMP_MAP_*'.  Also, the
    current special handling (incomplete, evidently) is the reason for the
PR104892
    misbehavior.  For PR100280 etc. we've resolved all such known ICEs --
removing
    the special handling for 'GOMP_MAP_FORCE_TOFROM' now resolves PR104892.

            PR middle-end/100280
            PR middle-end/104892
            gcc/
            * omp-oacc-kernels-decompose.cc (omp_oacc_kernels_decompose_1):
            Remove special handling of 'GOMP_MAP_FORCE_TOFROM'.
            gcc/testsuite/
            * c-c++-common/goacc/kernels-decompose-2.c: Adjust.
            * c-c++-common/goacc/kernels-decompose-pr100400-1-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr100400-1-2.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr100400-1-3.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr100400-1-4.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-2.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise.
            * c-c++-common/goacc/kernels-decompose-pr104774-1.c: Likewise.
            * gfortran.dg/goacc/classify-kernels.f95: Likewise.
            * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise.
            libgomp/
            * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Adjust.
            * testsuite/libgomp.oacc-c-c++-common/default-1.c: Likewise.
            * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
            Likewise.
            * testsuite/libgomp.oacc-c-c++-common/kernels-reduction-1.c:
            Likewise.
            * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise.
            * testsuite/libgomp.oacc-fortran/asyncwait-1.f90: Likewise.
            * testsuite/libgomp.oacc-fortran/kernels-reduction-1.f90:
            Likewise.