diff options
| author | Andrew Stubbs <ams@codesourcery.com> | 2020-01-30 18:25:15 +0000 |
|---|---|---|
| committer | Andrew Stubbs <ams@codesourcery.com> | 2020-04-23 15:58:09 +0100 |
| commit | ee9fcee3ec3a124dc3947c73c264bbcda97198df (patch) | |
| tree | 8b2b61a178417df1ce52beda7a233ecdc57a2200 | |
| parent | 901f5289d9465d4c388ae288f850ad4f29e99a2c (diff) | |
| download | gcc-ee9fcee3ec3a124dc3947c73c264bbcda97198df.zip gcc-ee9fcee3ec3a124dc3947c73c264bbcda97198df.tar.gz gcc-ee9fcee3ec3a124dc3947c73c264bbcda97198df.tar.bz2 | |
OpenACC: Avoid ICE in type-cast 'async', 'wait' clauses
2020-04-23 Andrew Stubbs <ams@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
PR middle-end/93488
gcc/
* omp-expand.c (expand_omp_target): Use force_gimple_operand_gsi on
t_async and the wait arguments.
gcc/testsuite/
* c-c++-common/goacc/pr93488.c: New file.
Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
| -rw-r--r-- | gcc/ChangeLog | 8 | ||||
| -rw-r--r-- | gcc/omp-expand.c | 13 | ||||
| -rw-r--r-- | gcc/testsuite/ChangeLog | 7 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/pr93488.c | 22 |
4 files changed, 46 insertions, 4 deletions
diff --git a/gcc/ChangeLog b/gcc/ChangeLog index 041894e..4e427c0 100644 --- a/gcc/ChangeLog +++ b/gcc/ChangeLog @@ -1,3 +1,11 @@ +2020-04-23 Andrew Stubbs <ams@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + + PR middle-end/93488 + + * omp-expand.c (expand_omp_target): Use force_gimple_operand_gsi on + t_async and the wait arguments. + 2020-04-23 Richard Sandiford <richard.sandiford@arm.com> PR tree-optimization/94727 diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index a642ccc..da1f4c3 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/ChangeLog b/gcc/testsuite/ChangeLog index 2b28272..c0caffb 100644 --- a/gcc/testsuite/ChangeLog +++ b/gcc/testsuite/ChangeLog @@ -1,3 +1,10 @@ +2020-04-23 Andrew Stubbs <ams@codesourcery.com> + Thomas Schwinge <thomas@codesourcery.com> + + PR middle-end/93488 + + * c-c++-common/goacc/pr93488.c: New file. + 2020-04-23 Richard Sandiford <richard.sandiford@arm.com> PR tree-optimization/94727 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 0000000..6fddad9 --- /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) +} |