* [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).