From: Andrew Stubbs <ams@codesourcery.com>
To: Thomas Schwinge <thomas@codesourcery.com>
Cc: <gcc-patches@gcc.gnu.org>, Jakub Jelinek <jakub@redhat.com>
Subject: Re: [PR93488] [OpenACC] ICE in type-cast 'async', 'wait' clauses
Date: Fri, 31 Jan 2020 00:05:00 -0000 [thread overview]
Message-ID: <ab9035d9-5b92-7ace-a0f8-521ad9566d6e@codesourcery.com> (raw)
In-Reply-To: <87r1zimomb.fsf@euler.schwinge.homeip.net>
[-- 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)
+}
next prev parent reply other threads:[~2020-01-30 21:13 UTC|newest]
Thread overview: 7+ messages / expand[flat|nested] mbox.gz Atom feed top
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 [this message]
2020-04-21 9:54 ` Thomas Schwinge
2020-04-23 15:02 ` Andrew Stubbs
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=ab9035d9-5b92-7ace-a0f8-521ad9566d6e@codesourcery.com \
--to=ams@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=thomas@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
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).