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

Revert "[nvptx, libgomp] Update pr85381-{2,4}.c test-cases" [PR89713, PR94392]

Browse source 
In response to PR94392 commit 75efe9cb1f
"c/94392 - only enable -ffinite-loops for C++", this reverts PR89713
commit 00908992f2, as apparently now again
"empty oacc loops are" no longer "removed before expand".

	libgomp/
	PR tree-optimization/89713
	PR c/94392
	* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Again expect
	'bar.sync'.
	* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise.
This commit is contained in:
Thomas Schwinge 2020-04-03 10:08:59 +02:00
parent 4441ecedc3
commit 2b1e849b35
3 changed files with 31 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

8
libgomp/ChangeLog
View file

@ -1,3 +1,11 @@
2020-04-03 Thomas Schwinge <thomas@codesourcery.com>
PR tree-optimization/89713
PR c/94392
* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Again expect
'bar.sync'.
* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise.
2020-03-31 Tobias Burnus <tobias@codesourcery.com>
* target.c (GOMP_target_enter_exit_data): Handle PSET/MAP_POINTER.

20
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c
View file

@ -15,4 +15,22 @@ main (void)
return 0;
}
/* { dg-final { scan-assembler-times "bar.sync" 0 } } */
/* Todo: Boths bar.syncs can be removed.
Atm we generate this dead code inbetween forked and joining:
mov.u32 %r28, %ntid.y;
mov.u32 %r29, %tid.y;
add.u32 %r30, %r29, %r29;
setp.gt.s32 %r31, %r30, 19;
@%r31 bra $L2;
add.u32 %r25, %r28, %r28;
mov.u32 %r24, %r30;
$L3:
add.u32 %r24, %r24, %r25;
setp.le.s32 %r33, %r24, 19;
@%r33 bra $L3;
$L2:
so the loop is not recognized as empty loop (which we detect by seeing if
joining immediately follows forked). */
/* { dg-final { scan-assembler-times "bar.sync" 2 } } */

5
libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c
View file

@ -21,4 +21,7 @@ main (void)
return 0;
}
/* { dg-final { scan-assembler-times "bar.sync" 0 } } */
/* Atm, %ntid.y is broadcast from one loop to the next, so there are 2 bar.syncs
for that (the other two are there for the same reason as in pr85381-2.c).
Todo: Recompute %ntid.y instead of broadcasting it. */
/* { dg-final { scan-assembler-times "bar.sync" 4 } } */