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.
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.
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.
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] * ...
(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); ...
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.
(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?
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.
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.
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>
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'.
GCC 9 branch is being closed
GCC 10.4 is being released, retargeting bugs to GCC 10.5.
GCC 10 branch is being closed.