On 21/04/2020 10:54, Thomas Schwinge wrote:
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>.
I've committed the attached.
Andrew
OpenACC: Avoid ICE in type-cast 'async', 'wait' clauses
2020-04-23 Andrew Stubbs <a...@codesourcery.com>
Thomas Schwinge <tho...@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 <tho...@codesourcery.com>
diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index a642ccc9980..da1f4c39d18 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)
+}