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>

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 041894e83326b4476b4d862f63fddf5bb4a0cc5b..4e427c06cf6e00b019e38ef9cc889e3a3ad0945b 100644 (file)
--- 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 a642ccc998041fcd82d8e56460afc49491c7d1ef..da1f4c39d18edceaa5ee3979d65a6a0cb2e09005 100644 (file)
--- 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 2b282721319f0d12e3c9087e314499e296c78b55..c0caffbd05f5d5c561d686efa37ac5d7c91346dd 100644 (file)
--- 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 (file)
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)
+}