Bug 104717 - [11 Regression] ICE: verify_ssa failed (Error: type mismatch between an SSA_NAME and its symbol)
Summary: [11 Regression] ICE: verify_ssa failed (Error: type mismatch between an SSA_N...
Status: ASSIGNED
Alias: None
Product: gcc
Classification: Unclassified
Component: fortran (show other bugs)
Version: 12.0
: P2 normal
Target Milestone: 11.5
Assignee: Jakub Jelinek
URL:
Keywords: openacc
Depends on:
Blocks:
 
Reported: 2022-02-28 12:32 UTC by Arseny Solokha
Modified: 2023-07-07 10:42 UTC (History)
3 users (show)

See Also:
Host:
Target:
Build:
Known to work:
Known to fail:
Last reconfirmed: 2022-02-28 00:00:00


Attachments
gcc12-pr104717.patch (737 bytes, patch)
2022-04-20 11:02 UTC, Jakub Jelinek
Details | Diff

Note You need to log in before you can comment on or make changes to this bug.
Description Arseny Solokha 2022-02-28 12:32:25 UTC
gfortran 12.0.1 20220227 snapshot (g:d1574a9b820f17adb9004255e2018967e9be063b) ICEs when compiling the following testcase, extracted from libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-2.f90, w/ -O1 -fopenacc -fstack-arrays:

program main
  implicit none (type, external)
  integer :: j
  integer, allocatable :: A(:)

  A = [(3*j, j=1, 10)]
  call foo (A, size(A))
  deallocate (A)
contains
  subroutine foo (array, nn)
    integer :: i, nn
    integer :: array(nn)

    !$acc parallel copyout(array)
    array = [(-i, i = 1, nn)]
    !$acc end parallel
  end subroutine foo
end

% gfortran-12.0.1 -O1 -fopenacc -fstack-arrays -c qatodltq.f90
qatodltq.f90:8:16:

    8 |   deallocate (A)
      |                ^
Error: type mismatch between an SSA_NAME and its symbol
qatodltq.f90:8:16: Error: type mismatch between an SSA_NAME and its symbol
while verifying SSA_NAME A.14_21 in statement
# .MEM_20 = VDEF <.MEM_13>
    A.14_21 = __builtin_alloca_with_align (_19, 32);
during IPA pass: pta
qatodltq.f90:8:16: internal compiler error: verify_ssa failed
0x1199793 verify_ssa(bool, bool)
	/var/tmp/portage/sys-devel/gcc-12.0.1_p20220227/work/gcc-12-20220227/gcc/tree-ssa.cc:1211
0xe66da5 execute_function_todo
	/var/tmp/portage/sys-devel/gcc-12.0.1_p20220227/work/gcc-12-20220227/gcc/passes.cc:2091
0xe67182 do_per_function
	/var/tmp/portage/sys-devel/gcc-12.0.1_p20220227/work/gcc-12-20220227/gcc/passes.cc:1694
0xe67182 do_per_function
	/var/tmp/portage/sys-devel/gcc-12.0.1_p20220227/work/gcc-12-20220227/gcc/passes.cc:1684
0xe671dc execute_todo
	/var/tmp/portage/sys-devel/gcc-12.0.1_p20220227/work/gcc-12-20220227/gcc/passes.cc:2138

BTW, the error is emitted twice in a row.
Comment 1 Jakub Jelinek 2022-02-28 13:09:06 UTC
Started with r6-5810-g597a8ab9c6f57c416c2b0a734c2098fa0335e628
As -fopenacc has been introduced already in r5-6458-g41dbbb3789850dfea98dd8984f69806284f87b6e , I think this means it is a regression from GCC 5.x.
Comment 2 Thomas Schwinge 2022-02-28 13:47:58 UTC
Reproduced, thanks for the report.


The problem disappears when adding '-fno-ipa-pta' to '-O1 -fopenacc -fstack-arrays'.  Not yet analyzed the differences.


The problem does not reproduce with '-O1 -fopenmp -fipa-pta -fstack-arrays' and code changes as per:

    -    !$acc parallel copyout(array)
    +    !$omp target map(from:array)
         array = [(-i, i = 1, nn)]
    -    !$acc end parallel
    +    !$omp end target

..., so really points to something specific to OpenACC handling.
Comment 3 Tom de Vries 2022-03-01 08:38:43 UTC
The test that is failing, is:
...
760       if (SSA_NAME_VAR (ssa_name) != NULL_TREE
761           && TREE_TYPE (ssa_name) != TREE_TYPE (SSA_NAME_VAR (ssa_name)))
762         {
763           error ("type mismatch between an SSA_NAME and its symbol");
764           return true;
765         }
...

For ssa-name:
...
(gdb) call debug_generic_expr (name)
A.14_21
(gdb) call debug_tree (name)
 <ssa_name 0x7ffff6351e10
    type <pointer_type 0x7ffff673abd0
        type <array_type 0x7ffff673ab28 type <integer_type 0x7ffff655b5e8 integer(kind=4)>
            sizes-gimplified BLK size <var_decl 0x7ffff634f3f0 D.4328> unit-size <var_decl 0x7ffff634f480 D.4329>
            align:32 warn_if_not_align:0 symtab:0 alias-set -1 structural-equality domain <integer_type 0x7ffff673aa80>
            pointer_to_this <pointer_type 0x7ffff673abd0>>
        public unsigned DI
        size <integer_cst 0x7ffff6543c00 constant 64>
        unit-size <integer_cst 0x7ffff6543c18 constant 8>
        align:64 warn_if_not_align:0 symtab:0 alias-set -1 structural-equality>
    var <var_decl 0x7ffff6356c60 A.14>
    def_stmt A.14_21 = __builtin_alloca_with_align (_19, 32);
    version:21
    ptr-info 0x7ffff63459f0>
...

The mismatch seems to be:
...
(gdb) call debug_generic_expr (name.typed.type)
integer(kind=4)[0:D.4266] *
(gdb) call debug_generic_expr (name.ssa_name.var.typed.type)
integer(kind=4)[0:D.4433] *
...
Comment 4 Tom de Vries 2022-03-01 11:01:19 UTC
(In reply to Tom de Vries from comment #3)
> The mismatch seems to be:
> ...
> (gdb) call debug_generic_expr (name.typed.type)
> integer(kind=4)[0:D.4266] *
> (gdb) call debug_generic_expr (name.ssa_name.var.typed.type)
> integer(kind=4)[0:D.4433] *
> ...

Setting a breakpoint at:
...
Breakpoint 4, ipa_pta_execute () at /home/vries/oacc/trunk/source-gcc/gcc/tree-ssa-structalias.cc:8322
8322          gcc_assert (!node->clone_of);
...
and doing:
...
(gdb) watch -l node->decl.function_decl.f.gimple_df->ssa_names.m_vecdata[21].ssa_name.var.typed.type
...
the second time we hit the breakpoint for foo.0._omp_fn.0, I hit the moment that the type changes:
...
(gdb) c
Continuing.

Hardware watchpoint 5: -location node->decl.function_decl.f.gimple_df->ssa_names.m_vecdata[21].ssa_name.var.typed.type

Old value = (tree_node *) 0x7ffff673abd0
New value = (tree_node *) 0x7ffff635d9d8
...
with:
...
(gdb) call debug_generic_expr ((tree_node *) 0x7ffff673abd0)
integer(kind=4)[0:D.4266] *
(gdb) call debug_generic_expr ((tree_node *) 0x7ffff635d9d8)
integer(kind=4)[0:D.4433] *
...

Backtrace:
...
(gdb) bt
#0  copy_tree_body_r (tp=0x7ffff634e3d8, walk_subtrees=0x7fffffffd04c, data=0x7fffffffd390)
    at /home/vries/oacc/trunk/source-gcc/gcc/tree-inline.cc:1418
#1  0x00000000011d4664 in walk_tree_1 (tp=0x7ffff634e3d8, tp@entry=0x7fffffffd098, 
    func=func@entry=0xf18070 <copy_tree_body_r(tree_node**, int*, void*)>, 
    data=data@entry=0x7fffffffd390, pset=pset@entry=0x0, lh=lh@entry=0x0)
    at /home/vries/oacc/trunk/source-gcc/gcc/tree.cc:11070
#2  0x0000000000f13ac4 in remap_decls (decls=<optimized out>, 
    nonlocalized_list=nonlocalized_list@entry=0x7ffff6342688, id=id@entry=0x7fffffffd390)
    at /home/vries/oacc/trunk/source-gcc/gcc/tree-inline.cc:791
#3  0x0000000000f15f10 in remap_block (block=block@entry=0x7fffffffd128, 
    id=id@entry=0x7fffffffd390) at /home/vries/oacc/trunk/source-gcc/gcc/tree-inline.cc:822
#4  0x0000000000f16096 in remap_blocks (block=block@entry=0x7ffff6342060, 
    id=id@entry=0x7fffffffd390) at /home/vries/oacc/trunk/source-gcc/gcc/tree-inline.cc:841
#5  0x0000000000f160c3 in remap_blocks (block=<optimized out>, id=id@entry=0x7fffffffd390)
    at /home/vries/oacc/trunk/source-gcc/gcc/tree-inline.cc:844
#6  0x0000000000f21622 in expand_call_inline (bb=<optimized out>, bb@entry=0x7ffff6358a28, 
    stmt=0x7ffff635b990, id=id@entry=0x7fffffffd390, to_purge=to_purge@entry=0x7fffffffd370)
    at /home/vries/oacc/trunk/source-gcc/gcc/tree-inline.cc:5033
#7  0x0000000000f22aea in gimple_expand_calls_inline (to_purge=0x7fffffffd370, 
    id=0x7fffffffd390, bb=0x7ffff6358a28)
    at /home/vries/oacc/trunk/source-gcc/gcc/tree-inline.cc:5320
--Type <RET> for more, q to quit, c to continue without paging--
#8  optimize_inline_calls (fn=0x7ffff673b100)
    at /home/vries/oacc/trunk/source-gcc/gcc/tree-inline.cc:5492
#9  0x0000000000beaa7c in inline_transform (node=<optimized out>)
    at /home/vries/oacc/trunk/source-gcc/gcc/ipa-inline-transform.cc:790
#10 0x0000000000d7f761 in execute_one_ipa_transform_pass (do_not_collect=true, 
    ipa_pass=0x2f52700, node=0x7ffff6343330)
    at /home/vries/oacc/trunk/source-gcc/gcc/passes.cc:2335
#11 execute_all_ipa_transforms (do_not_collect=do_not_collect@entry=true)
    at /home/vries/oacc/trunk/source-gcc/gcc/passes.cc:2398
#12 0x000000000094f301 in cgraph_node::get_body (this=this@entry=0x7ffff6343330)
    at /home/vries/oacc/trunk/source-gcc/gcc/cgraph.cc:4044
#13 0x00000000010ed712 in ipa_pta_execute ()
    at /home/vries/oacc/trunk/source-gcc/gcc/tree-ssa-structalias.cc:8320
#14 (anonymous namespace)::pass_ipa_pta::execute (this=<optimized out>)
    at /home/vries/oacc/trunk/source-gcc/gcc/tree-ssa-structalias.cc:8744
#15 0x0000000000d7fb91 in execute_one_pass (pass=pass@entry=0x2f52b90)
    at /home/vries/oacc/trunk/source-gcc/gcc/passes.cc:2643
#16 0x0000000000d816b2 in execute_ipa_pass_list (pass=0x2f52b90)
    at /home/vries/oacc/trunk/source-gcc/gcc/passes.cc:3083
#17 0x000000000095e0cb in symbol_table::compile (this=this@entry=0x7ffff6547000)
    at /home/vries/oacc/trunk/source-gcc/gcc/cgraphunit.cc:2311
--Type <RET> for more, q to quit, c to continue without paging--
#18 0x0000000000962130 in symbol_table::compile (this=0x7ffff6547000)
    at /home/vries/oacc/trunk/source-gcc/gcc/cgraphunit.cc:2532
#19 symbol_table::finalize_compilation_unit (this=0x7ffff6547000)
    at /home/vries/oacc/trunk/source-gcc/gcc/cgraphunit.cc:2529
#20 0x0000000000e7b44e in compile_file ()
    at /home/vries/oacc/trunk/source-gcc/gcc/toplev.cc:479
#21 0x00000000006fcd1f in do_compile (no_backend=false)
    at /home/vries/oacc/trunk/source-gcc/gcc/toplev.cc:2168
#22 toplev::main (this=this@entry=0x7fffffffd86e, argc=<optimized out>, argc@entry=18, 
    argv=<optimized out>, argv@entry=0x7fffffffd978)
    at /home/vries/oacc/trunk/source-gcc/gcc/toplev.cc:2320
#23 0x00000000006ff20b in main (argc=18, argv=0x7fffffffd978)
    at /home/vries/oacc/trunk/source-gcc/gcc/main.cc:39
...

I guess we're seeing the effect of:
...
      if (TREE_CODE (*tp) != OMP_CLAUSE)
        TREE_TYPE (*tp) = remap_type (TREE_TYPE (*tp), id);
...
Comment 5 Tom de Vries 2022-03-01 13:39:46 UTC
At original:
...
void foo ()
  ...
  #pragma acc parallel
    ...
        integer(kind=4) A.3[0:D.4266];
...

At gimple:
...
void foo ()
  ...
  #pragma omp target oacc_parallel
    ...
      integer(kind=4)[0:D.4266] * A.14;
      integer(kind=4) A.3[0:D.4266] [value-expr: *A.14];
      ...
      A.14 = __builtin_alloca_with_align (D.4329, 32);
...

And at ompexp:
...
void foo.0._omp_fn.0 ()
  ...
  integer(kind=4) A.3[0:D.4266] [value-expr: *A.14];
  integer(kind=4)[0:D.4266] * A.14;
  ...
  A.14 = __builtin_alloca_with_align (D.4392, 32);
...
and:
...
void foo ()
  ...
  __builtin_GOACC_parallel_keyed (-1, foo.0._omp_fn.0, 4, &.omp_data_arr.20,
                                  &.omp_data_sizes.21, &.omp_data_kinds.22, 0);
...

However, somehow the A.3 remains part of the BLOCK_VARS of foo, so when ipa inline (activated by pta-ipa, which does node->get_body ()) inlines foo into main, it does a remap of block-var A.3, and then a remap of its value-expr A.14, and then a remap of its TREE_TYPE.  Which has the effect of changing the TREE_TYPE of A.14 in foo.0._omp_fn.0.
Comment 6 Tom de Vries 2022-03-01 15:28:48 UTC
(In reply to Tom de Vries from comment #5)
> However, somehow the A.3 remains part of the BLOCK_VARS of foo, so when ipa
> inline (activated by pta-ipa, which does node->get_body ()) inlines foo into
> main, it does a remap of block-var A.3, and then a remap of its value-expr
> A.14, and then a remap of its TREE_TYPE.  Which has the effect of changing
> the TREE_TYPE of A.14 in foo.0._omp_fn.0.

Thomas, Jakub, any idea what the expected behaviour is here?

Should A.3 be a part of the BLOCK_VARS of foo, or not?
Comment 7 Jakub Jelinek 2022-03-07 16:30:13 UTC
As A.3 is in the body of a construct that is outlined into a separate function, it should be moved to the *omp_fn* function with it.
If you look at simple
int
main ()
{
  #pragma omp parallel
  {
    int a = 25;
    ++a;
  }
}
you can see in the gimple dump two pairs of {}s instead of just once:
    #pragma omp parallel
      {
        {
          int a;

          a = 25;
          a = a + 1;
        }
      }
(ditto target, task, (host) teams etc.), while with acc parallel that isn't the case:
  #pragma omp target oacc_parallel map(from:(*array.12_5) [len: D.4276]) map(alloc:array [pointer assign, bias: 0]) firstprivate(ubound.0) firstprivate(nn)
    {
The reason for that for the OpenMP construct is exactly to ensure that move_sese_region_to_fn does the right thing with it.
Comment 8 Jakub Jelinek 2022-04-20 11:02:28 UTC
Created attachment 52838 [details]
gcc12-pr104717.patch

Untested fix.
From brief look at #pragma acc parallel from C and C++ FEs, it looks correct there, so it is just Fortran.
Comment 9 GCC Commits 2022-04-25 21:14:43 UTC
The master branch has been updated by Thomas Schwinge <tschwinge@gcc.gnu.org>:

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

commit r12-8252-gb2202431910e30d8505c94d1cb9341cac7080d10
Author: Jakub Jelinek <jakub@redhat.com>
Date:   Wed Apr 20 19:06:17 2022 +0200

    fortran: Fix up gfc_trans_oacc_construct [PR104717]
    
    So that move_sese_region_to_fn works properly, OpenMP/OpenACC constructs
    for which that function is invoked need an extra artificial BIND_EXPR
    around their body so that we move all variables of the bodies.
    
    The C/C++ FEs do that both for OpenMP constructs like OMP_PARALLEL, OMP_TASK
    or OMP_TARGET and for OpenACC constructs that behave similarly to
    OMP_TARGET, but the Fortran FE only does that for OpenMP constructs.
    
    The following patch does that for OpenACC constructs too.
    
            PR fortran/104717
            gcc/fortran/
            * trans-openmp.cc (gfc_trans_oacc_construct): Wrap construct body
            in an extra BIND_EXPR.
            gcc/testsuite/
            * gfortran.dg/goacc/pr104717.f90: New test.
            * gfortran.dg/goacc/privatization-1-compute-loop.f90: Adjust.
            libgomp/
            * testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Adjust.
    
    Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
Comment 10 GCC Commits 2022-04-28 13:21:50 UTC
The master branch has been updated by Thomas Schwinge <tschwinge@gcc.gnu.org>:

https://gcc.gnu.org/g:2a570f11a2fecf23998d7fe1d5cabad62cfe5cec

commit r12-8311-g2a570f11a2fecf23998d7fe1d5cabad62cfe5cec
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Tue Apr 26 18:41:23 2022 +0200

    Fix up 'libgomp.oacc-fortran/print-1.f90' GCN offloading compilation [PR104717]
    
    That got broken by recent commit b2202431910e30d8505c94d1cb9341cac7080d10
    "fortran: Fix up gfc_trans_oacc_construct [PR104717]".
    
            PR fortran/104717
            libgomp/
            * testsuite/libgomp.oacc-fortran/print-1.f90: Add OpenACC
            privatization scanning.  For GCN offloading compilation, raise
            '-mgang-private-size'.
Comment 11 Richard Biener 2022-05-27 09:47:38 UTC
GCC 9 branch is being closed
Comment 12 Jakub Jelinek 2022-06-28 10:48:33 UTC
GCC 10.4 is being released, retargeting bugs to GCC 10.5.
Comment 13 Richard Biener 2023-07-07 10:42:38 UTC
GCC 10 branch is being closed.