[openacc] Fix ICE verifying gimple

Message ID d5b866af-3781-abe5-9561-a33ac32f4913@codesourcery.com
State New
Headers show
Series
  • [openacc] Fix ICE verifying gimple
Related show

Commit Message

Andrew Stubbs Nov. 22, 2019, 11:06 a.m.
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

Comments

Andrew Stubbs Jan. 16, 2020, 1:46 p.m. | #1
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?

>
Thomas Schwinge Jan. 29, 2020, 12:30 p.m. | #2
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 ();
Andrew Stubbs Jan. 29, 2020, 12:54 p.m. | #3
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
Andrew Stubbs Jan. 30, 2020, 9:12 p.m. | #4
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
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)
+}

Patch

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 ();