Here's a fix for another flaky test.
Tested on x86_64-linux-gnu by running it with gfx908 and
gfx90a accelerators, 500 times for each of those.
OK for trunk?
TIA
---------- >8 ----------
In r11-3059-g8183ebcdc1c843, Julian fixed a few issues with
atomic_capture-2.c relying on iteration order guarantees that do not
exist under OpenACC parallelized loops and, notably, do not happen even
by accident on AMDGCN.
The atomic_capture-3.c testcase was made by copying it from
atomic_capture-2.c and adding additional options in commit
r12-310-g4cf3b10f27b199, but from an older version of
atomic_capture-2.c, which lacked these ordering fixes fixes, so they
resurfaced in this test.
This patch ports those fixes from atomic_capture-2.c into
atomic_capture-3.c.
libgomp/ChangeLog:
* testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c: Copy
changes in r11-3059-g8183ebcdc1c843 from atomic_capture-2.c.
---
.../atomic_capture-3.c | 92 +++++++++----------
1 file changed, 43 insertions(+), 49 deletions(-)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c
index b976094998f2..b8a76a17560d 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c
@@ -38,11 +38,9 @@ main(int argc, char **argv)
imin = idata[i] < imin ? idata[i] : imin;
}
- if (imax != 1234 || imin != 0)
+ if (imax != 1234 || imin < 0 || imin > 1)
abort ();
- return 0;
-
igot = 0;
iexp = 32;
@@ -444,17 +442,16 @@ main(int argc, char **argv)
}
}
+ int ones = 0, zeros = 0;
+
for (i = 0; i < N; i++)
- if (i % 2 == 0)
- {
- if (idata[i] != 1)
- abort ();
- }
- else
- {
- if (idata[i] != 0)
- abort ();
- }
+ if (idata[i] == 1)
+ ones++;
+ else if (idata[i] == 0)
+ zeros++;
+
+ if (ones != N / 2 || zeros != N / 2)
+ abort ();
if (iexp != igot)
abort ();
@@ -492,17 +489,16 @@ main(int argc, char **argv)
}
}
+ ones = zeros = 0;
+
for (i = 0; i < N; i++)
- if (i % 2 == 0)
- {
- if (idata[i] != 0)
- abort ();
- }
- else
- {
- if (idata[i] != 1)
- abort ();
- }
+ if (idata[i] == 1)
+ ones++;
+ else if (idata[i] == 0)
+ zeros++;
+
+ if (ones != N / 2 || zeros != N / 2)
+ abort ();
if (iexp != igot)
abort ();
@@ -580,7 +576,7 @@ main(int argc, char **argv)
if (lexp != lgot)
abort ();
- lgot = 2LL;
+ lgot = 2LL << N;
lexp = 2LL;
#pragma acc data copy (lgot, ldata[0:N])
@@ -588,7 +584,7 @@ main(int argc, char **argv)
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
- long long expr = 1LL << N;
+ long long expr = 2LL;
#pragma acc atomic capture
{ lgot = lgot / expr; ldata[i] = lgot; }
@@ -1451,17 +1447,16 @@ main(int argc, char **argv)
}
}
+ ones = zeros = 0;
+
for (i = 0; i < N; i++)
- if (i % 2 == 0)
- {
- if (fdata[i] != 1.0)
- abort ();
- }
- else
- {
- if (fdata[i] != 0.0)
- abort ();
- }
+ if (fdata[i] == 1.0)
+ ones++;
+ else if (fdata[i] == 0.0)
+ zeros++;
+
+ if (ones != N / 2 || zeros != N / 2)
+ abort ();
if (fexp != fgot)
abort ();
@@ -1499,17 +1494,16 @@ main(int argc, char **argv)
}
}
+ ones = zeros = 0;
+
for (i = 0; i < N; i++)
- if (i % 2 == 0)
- {
- if (fdata[i] != 0.0)
- abort ();
- }
- else
- {
- if (fdata[i] != 1.0)
- abort ();
- }
+ if (fdata[i] == 1.0)
+ ones++;
+ else if (fdata[i] == 0.0)
+ zeros++;
+
+ if (ones != N / 2 || zeros != N / 2)
+ abort ();
if (fexp != fgot)
abort ();
@@ -1570,7 +1564,7 @@ main(int argc, char **argv)
abort ();
fgot = 8192.0*8192.0*64.0;
- fexp = 1.0;
+ fexp = fgot;
#pragma acc data copy (fgot, fdata[0:N])
{
@@ -1587,15 +1581,15 @@ main(int argc, char **argv)
if (fexp != fgot)
abort ();
- fgot = 4.0;
- fexp = 4.0;
+ fgot = 2.0 * (1LL << N);
+ fexp = 2.0;
#pragma acc data copy (fgot, fdata[0:N])
{
#pragma acc parallel loop
for (i = 0; i < N; i++)
{
- long long expr = 1LL << N;
+ long long expr = 2LL;
#pragma acc atomic capture
{ fgot = fgot / expr; fdata[i] = fgot; }
--
2.52.0
