[patch, openacc] Fix ICE verifying gimple

[patch, openacc] Fix ICE verifying gimple

[Andrew Stubbs]

[-- Attachment #1: Type: text/plain, Size: 1013 bytes --]

This test case causes an ICE (reformatted for email):

   void test(int k)
   {
     unsigned int x = 1;
   #pragma acc parallel loop async(x)
     for (int i = 0; i < k; i++) { }
   }

   t.c: In function 'test':
   t.c:4:9: error: invalid argument to gimple call
       4 | #pragma acc parallel loop async(x)
         |         ^~~
   (int) x
   __builtin_GOACC_parallel_keyed (-1, test._omp_fn.0, 1,
                                   &.omp_data_arr.3, &.omp_data_sizes.4,
                                   &.omp_data_kinds.5, 536936447,
                                  (int) x, 0);
   during GIMPLE pass: ompexp
   dump file: t.c.013t.ompexp
   t.c:4:9: internal compiler error: verify_gimple failed

The problem is that "x" needs to be cast to "int" (from "unsigned int") 
before calling the function, and that's not valid in a gimple call.

The attached patch assigns the "(int) x" to a temporary and passes that 
to the function instead.

OK to commit?

-- 
Andrew Stubbs
CodeSourcery / Mentor Graphics

[-- Attachment #2: 191122-fix-goacc-parallel-keyed.patch --]
[-- Type: text/x-patch, Size: 779 bytes --]

Normalize GOACC_parallel_keyed async parameter.

2019-11-22  Andrew Stubbs  <ams@codesourcery.com>

	gcc/
	* omp-expand.c (expand_omp_target): Pass sync parameter to
	GOACC_parallel_keyed via a temporary variable.

diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 6f945011cf5..08f95587e95 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -8418,7 +8418,12 @@ expand_omp_target (struct omp_region *region)
 					      i_async));
 	  }
 	if (t_async)
-	  args.safe_push (t_async);
+	  {
+	    tree tmp = create_tmp_var (TREE_TYPE (t_async));
+	    gimple *stmt = gimple_build_assign (tmp, t_async);
+	    gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
+	    args.safe_push (tmp);
+	  }
 
 	/* Save the argument index, and ... */
 	unsigned t_wait_idx = args.length ();

Re: [patch, openacc] Fix ICE verifying gimple

[Andrew Stubbs]

Ping.

On 22/11/2019 11:06, Andrew Stubbs wrote:
> This test case causes an ICE (reformatted for email):
> 
>    void test(int k)
>    {
>      unsigned int x = 1;
>    #pragma acc parallel loop async(x)
>      for (int i = 0; i < k; i++) { }
>    }
> 
>    t.c: In function 'test':
>    t.c:4:9: error: invalid argument to gimple call
>        4 | #pragma acc parallel loop async(x)
>          |         ^~~
>    (int) x
>    __builtin_GOACC_parallel_keyed (-1, test._omp_fn.0, 1,
>                                    &.omp_data_arr.3, &.omp_data_sizes.4,
>                                    &.omp_data_kinds.5, 536936447,
>                                   (int) x, 0);
>    during GIMPLE pass: ompexp
>    dump file: t.c.013t.ompexp
>    t.c:4:9: internal compiler error: verify_gimple failed
> 
> The problem is that "x" needs to be cast to "int" (from "unsigned int") 
> before calling the function, and that's not valid in a gimple call.
> 
> The attached patch assigns the "(int) x" to a temporary and passes that 
> to the function instead.
> 
> OK to commit?
> 

[PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses (was: [patch, openacc] Fix ICE verifying gimple)

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 3114 bytes --]

Hi Andrew!

On 2019-11-22T11:06:14+0000, Andrew Stubbs <ams@codesourcery.com> wrote:
> This test case causes an ICE (reformatted for email):
>
>    void test(int k)
>    {
>      unsigned int x = 1;
>    #pragma acc parallel loop async(x)
>      for (int i = 0; i < k; i++) { }
>    }

Curious, why didn't you include such a test case in your actual patch?

>    t.c: In function 'test':
>    t.c:4:9: error: invalid argument to gimple call
>        4 | #pragma acc parallel loop async(x)
>          |         ^~~
>    (int) x
>    __builtin_GOACC_parallel_keyed (-1, test._omp_fn.0, 1,
>                                    &.omp_data_arr.3, &.omp_data_sizes.4,
>                                    &.omp_data_kinds.5, 536936447,
>                                   (int) x, 0);
>    during GIMPLE pass: ompexp
>    dump file: t.c.013t.ompexp
>    t.c:4:9: internal compiler error: verify_gimple failed
>
> The problem is that "x" needs to be cast to "int" (from "unsigned int") 
> before calling the function, and that's not valid in a gimple call.

ACK, thanks.

> The attached patch assigns the "(int) x" to a temporary and passes that 
> to the function instead.
>
> OK to commit?

A little bit more thoroughness, please.  :-)

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.


Grüße
 Thomas


> Normalize GOACC_parallel_keyed async parameter.
>
> 2019-11-22  Andrew Stubbs  <ams@codesourcery.com>
>
> 	gcc/
> 	* omp-expand.c (expand_omp_target): Pass sync parameter to
> 	GOACC_parallel_keyed via a temporary variable.
>
> diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
> index 6f945011cf5..08f95587e95 100644
> --- a/gcc/omp-expand.c
> +++ b/gcc/omp-expand.c
> @@ -8418,7 +8418,12 @@ expand_omp_target (struct omp_region *region)
>  					      i_async));
>  	  }
>  	if (t_async)
> -	  args.safe_push (t_async);
> +	  {
> +	    tree tmp = create_tmp_var (TREE_TYPE (t_async));
> +	    gimple *stmt = gimple_build_assign (tmp, t_async);
> +	    gsi_insert_before (&gsi, stmt, GSI_SAME_STMT);
> +	    args.safe_push (tmp);
> +	  }
>  
>  	/* Save the argument index, and ... */
>  	unsigned t_wait_idx = args.length ();

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]

Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses

[Andrew Stubbs]

On 29/01/2020 12:30, Thomas Schwinge wrote:
> Hi Andrew!
> 
> On 2019-11-22T11:06:14+0000, Andrew Stubbs <ams@codesourcery.com> wrote:
>> This test case causes an ICE (reformatted for email):
>>
>>     void test(int k)
>>     {
>>       unsigned int x = 1;
>>     #pragma acc parallel loop async(x)
>>       for (int i = 0; i < k; i++) { }
>>     }
> 
> Curious, why didn't you include such a test case in your actual patch?

Because I'm a terrible person.

>>     t.c: In function 'test':
>>     t.c:4:9: error: invalid argument to gimple call
>>         4 | #pragma acc parallel loop async(x)
>>           |         ^~~
>>     (int) x
>>     __builtin_GOACC_parallel_keyed (-1, test._omp_fn.0, 1,
>>                                     &.omp_data_arr.3, &.omp_data_sizes.4,
>>                                     &.omp_data_kinds.5, 536936447,
>>                                    (int) x, 0);
>>     during GIMPLE pass: ompexp
>>     dump file: t.c.013t.ompexp
>>     t.c:4:9: internal compiler error: verify_gimple failed
>>
>> The problem is that "x" needs to be cast to "int" (from "unsigned int")
>> before calling the function, and that's not valid in a gimple call.
> 
> ACK, thanks.
> 
>> The attached patch assigns the "(int) x" to a temporary and passes that
>> to the function instead.
>>
>> OK to commit?
> 
> A little bit more thoroughness, please.  :-)
> 
> 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.

OK, I'll take a look at your testcase and cook up a patch.

Today I learned about force_gimple_operand_gsi. :-)

Andrew

Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses

[Andrew Stubbs]

[-- Attachment #1: Type: text/plain, Size: 1042 bytes --]

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?

Andrew

[-- Attachment #2: 200130-normalize-goacc-args.patch --]
[-- Type: text/x-patch, Size: 2152 bytes --]

Normalize GOACC_parallel_keyed async and wait parameters

2020-01-30  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.

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)
+}

Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses

[Thomas Schwinge]

Hi Andrew!

On 2020-01-30T21:12:59+0000, Andrew Stubbs <ams@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
<thomas@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

Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses

[Andrew Stubbs]

[-- Attachment #1: Type: text/plain, Size: 473 bytes --]

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
> <thomas@codesourcery.com>" in the commit log, see
> <https://gcc.gnu.org/wiki/Reviewed-by>.

I've committed the attached.

Andrew

[-- Attachment #2: 200423-async-wait-fix.patch --]
[-- Type: text/x-patch, Size: 2267 bytes --]

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/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)
+}