https://gcc.gnu.org/bugzilla/show_bug.cgi?id=102330
Thomas Schwinge <tschwinge at gcc dot gnu.org> changed:
What |Removed |Added
----------------------------------------------------------------------------
CC| |rguenth at gcc dot gnu.org
--- Comment #8 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Posting my current status here, in case anybody feels like verifying my
thoughts. :-)
So I'm looking into Jakub's last, simplest reproducer:
(In reply to Jakub Jelinek from comment #7)
> program p
> i = 0
> !$omp task shared(i)
> i = 1
> !$omp end task
> !$omp taskwait
> !$acc parallel loop
> do i = 1, 8
> end do
> end
>
> also ICEs
ACK.
> it just needs the iterator that is marked addressable during
> omp-low.c which has code to regimplify statements that access such
> variables, but it seems the problematic stmts are emitted only during
> omp-expand.c.
ACK, 'gcc/omp-expand.cc:expand_oacc_for'.
> The various expand_omp_* loop expansions have DECL_P (...) &&
> TREE_ADDRESSABLE (...) conditions in various spots, but perhaps OpenACC loop
> expansion doesn't.
(Thanks, that has help me with PR104132!)
However, that doesn't seem to be the problem here: we get reported
"non-register as LHS of binary operation", but we've got in
'gcc/omp-expand.cc:expand_oacc_for':
7868 ass = gimple_build_assign (v, expr);
(gdb) n
7869 gsi_insert_before (&gsi, ass, GSI_SAME_STMT);
(gdb) call debug_gimple_stmt(ass)
i = 1 + .offset.8;
(gdb) call debug_tree(v)
<var_decl 0x7ffff7fc3e10 i
type <integer_type 0x7ffff75555e8 integer(kind=4) public SI
size <integer_cst 0x7ffff753ee40 constant 32>
unit-size <integer_cst 0x7ffff753ee58 constant 4>
align:32 warn_if_not_align:0 symtab:0 alias-set -1 canonical-type
0x7ffff75555e8 precision:32 min <integer_cst 0x7ffff753edf8 -2147483648> max
<integer_cst 0x7ffff753ee10 2147483647>
pointer_to_this <pointer_type 0x7ffff755da80>>
used SI
source-gcc/gcc/testsuite/gfortran.dg/goacc-gomp/pr102330-3.f90:9:3 size
<integer_cst 0x7ffff753ee40 32> unit-size <integer_cst 0x7ffff753ee58 4>
align:32 warn_if_not_align:0 context <function_decl 0x7ffff7730d00 p>>
(gdb) call is_gimple_lvalue(v)
$4 = true
(gdb) call is_gimple_reg(v)
$5 = true
(gdb) print v.base.addressable_flag
$6 = 0
So, 'i' is not 'TREE_ADDRESSABLE'.
Is the problem rather, that we earlier have done wrong, as Richard pointed out:
(In reply to Richard Biener from comment #1)
> Confirmed. omp expansion seems to introudce this non-gimple code and likely
> also makes 'i' not a register (for OACC):
>
> .data_dep.5D.4044 = .UNIQUE (OACC_PRIVATE, .data_dep.5D.4044, -1, &iD.4045);
That's built in 'gcc/omp-low.cc:lower_oacc_private_marker'. It's not marked
'TREE_ADDRESSABLE' (per debug log above), but the address of 'i' has been taken
(here).
> but
>
> iD.4045 = 1 + .offset.9D.4049;
>
> and
>
> iD.4045 = 1 + 8;
>
> (which is also unfolded)
Assuming 'TREE_ADDRESSABLE', that last item is resolved via my proposed
PR104132 change:
[...]
+gimple_simplified to D.4257 = 9;
[...]
- iD.4257 = 1 + 8;
+ D.4266 = 9;
+ iD.4259 = D.4266;
[...]
So if in 'gcc/omp-expand.cc:expand_oacc_for' I force 'TREE_ADDRESSABLE', then
the ICE goes away, and we get valid GIMPLE generated, per my understanding.
So, something like this maybe (untested)?
--- gcc/omp-low.cc
+++ gcc/omp-low.cc
@@ -11513,6 +11513,7 @@ lower_oacc_private_marker (omp_context *ctx)
}
gcc_checking_assert (decl);
+ TREE_ADDRESSABLE (decl) = 1;
tree addr = build_fold_addr_expr (decl);
args.safe_push (addr);
}
However: given that the '#pragma acc loop' has an implicit 'private(i)' -- why
does the 'gcc/omp-expand.cc:expand_oacc_for' even bother about the "outer 'i'"?
The loop variable 'i' should be a completely separate, local, register-like
temporary. Amy I misunderstanding something, or what's going wrong, elsewhere?
