diff options
author | Julian Brown <julian@codesourcery.com> | 2020-09-07 11:43:16 -0700 |
---|---|---|
committer | Julian Brown <julian@codesourcery.com> | 2020-09-08 13:26:42 -0700 |
commit | 8183ebcdc1c843f15c807e5bc26dbafe4e8c4dc3 (patch) | |
tree | 076fc447aecedd26bf0fb5252c2caa5bc36ba3ef /libgomp | |
parent | e929d65b48ad5583e995d872ca5de95b23de4d72 (diff) | |
download | gcc-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.c | 92 |
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; } |