Hi Andrew! On 2020-01-30T21:12:59+0000, Andrew Stubbs <a...@codesourcery.com> wrote: > On 29/01/2020 12:30, Thomas Schwinge wrote: >> As can be seen in the code a few lines below, the very same problem also >> exists for the 'wait' clause; it seems reasonable to fix both at the same >> time. This is not a recent regression, but a user-visible valid-code ICE >> that has existed "forever"; I filed <https://gcc.gnu.org/PR93488> "ICE in >> type-cast 'async', 'wait' clauses" for tracking. This problem is similar >> to the OpenMP 'device' clause <https://gcc.gnu.org/PR71758> "ICE in >> verify_gimple_in_cfg, at tree-cfg.c:5212"; I suggest we also use >> 'force_gimple_operand_gsi' instead of manually doing your suggested >> 'create_tmp_var', 'gimple_build_assign', 'gsi_insert_before'. Include a >> test case that covers all relevant code paths; I've attached a test case >> to the PR, but I've not verified whether "that covers *all* relevant code >> paths". This should then be backported to all GCC release branches; I >> can easily test the backports for you, if you're not already set up to do >> such testing. > > How's this?
Thanks, looks good! > Normalize GOACC_parallel_keyed async and wait parameters Oh, not only 'GOACC_parallel_keyed', but also other OpenACC directives. Maybe simply "[OpenACC] Avoid ICE in type-cast 'async', 'wait' clauses"? ;-P To record the review effort, please include "Reviewed-by: Thomas Schwinge <tho...@codesourcery.com>" in the commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>. Grüße Thomas > diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c > index cd423ad799e..ec4baf46965 100644 > --- a/gcc/omp-expand.c > +++ b/gcc/omp-expand.c > @@ -8418,7 +8418,9 @@ expand_omp_target (struct omp_region *region) > i_async)); > } > if (t_async) > - args.safe_push (t_async); > + args.safe_push (force_gimple_operand_gsi (&gsi, t_async, true, > + NULL_TREE, true, > + GSI_SAME_STMT)); > > /* Save the argument index, and ... */ > unsigned t_wait_idx = args.length (); > @@ -8431,9 +8433,12 @@ expand_omp_target (struct omp_region *region) > for (; c; c = OMP_CLAUSE_CHAIN (c)) > if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_WAIT) > { > - args.safe_push (fold_convert_loc (OMP_CLAUSE_LOCATION (c), > - integer_type_node, > - OMP_CLAUSE_WAIT_EXPR (c))); > + tree arg = fold_convert_loc (OMP_CLAUSE_LOCATION (c), > + integer_type_node, > + OMP_CLAUSE_WAIT_EXPR (c)); > + arg = force_gimple_operand_gsi (&gsi, arg, true, NULL_TREE, true, > + GSI_SAME_STMT); > + args.safe_push (arg); > num_waits++; > } > > diff --git a/gcc/testsuite/c-c++-common/goacc/pr93488.c > b/gcc/testsuite/c-c++-common/goacc/pr93488.c > new file mode 100644 > index 00000000000..6fddad919d2 > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/goacc/pr93488.c > @@ -0,0 +1,22 @@ > +/* PR middle-end/93488 > + > + Ensure that wait and async arguments can be cast to the correct type > + without breaking gimple verification. */ > + > +void test() > +{ > + /* int */ unsigned char a = 1; > + /* int */ unsigned char w = 1; > + > +#pragma acc parallel wait(w) async(a) > + ; > +#pragma acc kernels wait(w) async(a) > + ; > +#pragma acc serial wait(w) async(a) > + ; > + int data = 0; > +#pragma acc enter data wait(w) async(a) create(data) > +#pragma acc update wait(w) async(a) device(data) > +#pragma acc exit data wait(w) async(a) delete(data) > +#pragma acc wait(w) async(a) > +} ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
