diff options
| author | Thomas Schwinge <thomas@codesourcery.com> | 2020-04-03 10:08:59 +0200 |
|---|---|---|
| committer | Thomas Schwinge <thomas@codesourcery.com> | 2020-04-03 10:10:07 +0200 |
| commit | 2b1e849b35bfe694b09f090397944e11c9f1dc04 (patch) | |
| tree | bb6cae66a8ab0124ce9e3471cdfd0ba8fc3dbd63 /libgomp | |
| parent | 4441ecedc3d74114683761901075085651928c52 (diff) | |
Revert "[nvptx, libgomp] Update pr85381-{2,4}.c test-cases" [PR89713, PR94392]
In response to PR94392 commit 75efe9cb1f8938a713ce540dc3b27bc2afcd3fae
"c/94392 - only enable -ffinite-loops for C++", this reverts PR89713
commit 00908992f2a78f213d227aea8dbab014a1361df0, 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.
Diffstat (limited to 'libgomp')
| -rw-r--r-- | libgomp/ChangeLog | 8 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c | 20 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c | 5 |
3 files changed, 31 insertions, 2 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 6c437930b02..3716f559aa1 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -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. diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c index 2cb5b95949d..6570c64afff 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-2.c @@ -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 } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c index e8a433ffc0a..d955d79718d 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr85381-4.c @@ -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 } } */ |