* [patch, openacc] Fix ICE verifying gimple
@ 2019-11-22 11:18 Andrew Stubbs
2020-01-16 14:44 ` Andrew Stubbs
2020-01-29 12:42 ` [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses (was: [patch, openacc] Fix ICE verifying gimple) Thomas Schwinge
0 siblings, 2 replies; 7+ messages in thread
From: Andrew Stubbs @ 2019-11-22 11:18 UTC (permalink / raw)
To: gcc-patches
[-- 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 ();
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [patch, openacc] Fix ICE verifying gimple
2019-11-22 11:18 [patch, openacc] Fix ICE verifying gimple Andrew Stubbs
@ 2020-01-16 14:44 ` Andrew Stubbs
2020-01-29 12:42 ` [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses (was: [patch, openacc] Fix ICE verifying gimple) Thomas Schwinge
1 sibling, 0 replies; 7+ messages in thread
From: Andrew Stubbs @ 2020-01-16 14:44 UTC (permalink / raw)
To: gcc-patches, Thomas Schwinge
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?
>
^ permalink raw reply [flat|nested] 7+ messages in thread
* [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses (was: [patch, openacc] Fix ICE verifying gimple)
2019-11-22 11:18 [patch, openacc] Fix ICE verifying gimple Andrew Stubbs
2020-01-16 14:44 ` Andrew Stubbs
@ 2020-01-29 12:42 ` Thomas Schwinge
2020-01-29 13:32 ` [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses Andrew Stubbs
2020-01-31 0:05 ` Andrew Stubbs
1 sibling, 2 replies; 7+ messages in thread
From: Thomas Schwinge @ 2020-01-29 12:42 UTC (permalink / raw)
To: Andrew Stubbs; +Cc: gcc-patches, Jakub Jelinek
[-- 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 --]
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses
2020-01-29 12:42 ` [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses (was: [patch, openacc] Fix ICE verifying gimple) Thomas Schwinge
@ 2020-01-29 13:32 ` Andrew Stubbs
2020-01-31 0:05 ` Andrew Stubbs
1 sibling, 0 replies; 7+ messages in thread
From: Andrew Stubbs @ 2020-01-29 13:32 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Jakub Jelinek
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
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses
2020-01-29 12:42 ` [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses (was: [patch, openacc] Fix ICE verifying gimple) Thomas Schwinge
2020-01-29 13:32 ` [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses Andrew Stubbs
@ 2020-01-31 0:05 ` Andrew Stubbs
2020-04-21 9:54 ` Thomas Schwinge
1 sibling, 1 reply; 7+ messages in thread
From: Andrew Stubbs @ 2020-01-31 0:05 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches, Jakub Jelinek
[-- 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)
+}
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses
2020-01-31 0:05 ` Andrew Stubbs
@ 2020-04-21 9:54 ` Thomas Schwinge
2020-04-23 15:02 ` Andrew Stubbs
0 siblings, 1 reply; 7+ messages in thread
From: Thomas Schwinge @ 2020-04-21 9:54 UTC (permalink / raw)
To: Andrew Stubbs; +Cc: gcc-patches, Jakub Jelinek
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
^ permalink raw reply [flat|nested] 7+ messages in thread
* Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses
2020-04-21 9:54 ` Thomas Schwinge
@ 2020-04-23 15:02 ` Andrew Stubbs
0 siblings, 0 replies; 7+ messages in thread
From: Andrew Stubbs @ 2020-04-23 15:02 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: Jakub Jelinek, gcc-patches
[-- 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)
+}
^ permalink raw reply [flat|nested] 7+ messages in thread
end of thread, other threads:[~2020-04-23 15:02 UTC | newest]
Thread overview: 7+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2019-11-22 11:18 [patch, openacc] Fix ICE verifying gimple Andrew Stubbs
2020-01-16 14:44 ` Andrew Stubbs
2020-01-29 12:42 ` [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses (was: [patch, openacc] Fix ICE verifying gimple) Thomas Schwinge
2020-01-29 13:32 ` [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses Andrew Stubbs
2020-01-31 0:05 ` Andrew Stubbs
2020-04-21 9:54 ` Thomas Schwinge
2020-04-23 15:02 ` Andrew Stubbs
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).