aboutsummaryrefslogtreecommitdiff
path: root/libgomp
diff options
context:
space:
mode:
authorJulian Brown <julian@codesourcery.com>2020-09-07 11:43:16 -0700
committerJulian Brown <julian@codesourcery.com>2020-09-08 13:26:42 -0700
commit8183ebcdc1c843f15c807e5bc26dbafe4e8c4dc3 (patch)
tree076fc447aecedd26bf0fb5252c2caa5bc36ba3ef /libgomp
parente929d65b48ad5583e995d872ca5de95b23de4d72 (diff)
downloadgcc-8183ebcdc1c843f15c807e5bc26dbafe4e8c4dc3.zip
gcc-8183ebcdc1c843f15c807e5bc26dbafe4e8c4dc3.tar.gz
gcc-8183ebcdc1c843f15c807e5bc26dbafe4e8c4dc3.tar.bz2
openacc: Fix atomic_capture-2.c iteration-ordering issues
The test case was written with assumptions about loop iteration ordering that are not guaranteed by OpenACC and do not apply on all targets, in particular AMD GCN. This patch removes those assumptions. 2020-09-08 Julian Brown <julian@codesourcery.com> libgomp/ * testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c: Remove iteration-ordering assumptions.
Diffstat (limited to 'libgomp')
-rw-r--r--libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c92
1 files changed, 43 insertions, 49 deletions
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c
index 842f2de..4f83f03 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c
@@ -37,11 +37,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;
@@ -443,17 +441,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 ();
@@ -491,17 +488,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 ();
@@ -579,7 +575,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])
@@ -587,7 +583,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; }
@@ -1450,17 +1446,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 ();
@@ -1498,17 +1493,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 ();
@@ -1569,7 +1563,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])
{
@@ -1586,15 +1580,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; }