procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

OpenACC: Avoid ICE in type-cast 'async', 'wait' clauses

Browse source 
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>
This commit is contained in:
Andrew Stubbs 2020-01-30 18:25:15 +00:00
parent 901f5289d9
commit ee9fcee3ec
4 changed files with 46 additions and 4 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

8
gcc/ChangeLog
View file

@ -1,3 +1,11 @@
2020-04-23 Andrew Stubbs <ams@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
PR middle-end/93488
* omp-expand.c (expand_omp_target): Use force_gimple_operand_gsi on
t_async and the wait arguments.
2020-04-23 Richard Sandiford <richard.sandiford@arm.com>
PR tree-optimization/94727

13
gcc/omp-expand.c
View file

@ -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++;
}

7
gcc/testsuite/ChangeLog
View file

@ -1,3 +1,10 @@
2020-04-23 Andrew Stubbs <ams@codesourcery.com>
Thomas Schwinge <thomas@codesourcery.com>
PR middle-end/93488
* c-c++-common/goacc/pr93488.c: New file.
2020-04-23 Richard Sandiford <richard.sandiford@arm.com>
PR tree-optimization/94727

22
gcc/testsuite/c-c++-common/goacc/pr93488.c Normal file
View file

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