From: Thomas Schwinge Date: Fri, 3 Apr 2020 08:08:59 +0000 (+0200) Subject: Revert "[nvptx, libgomp] Update pr85381-{2,4}.c test-cases" [PR89713, PR94392] X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=2b1e849b35bfe694b09f090397944e11c9f1dc04;p=gcc.git 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. --- 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 + + 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 * 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 } } */