From: Julian Brown Date: Mon, 7 Sep 2020 18:43:16 +0000 (-0700) Subject: openacc: Fix atomic_capture-2.c iteration-ordering issues X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=8183ebcdc1c843f15c807e5bc26dbafe4e8c4dc3;p=gcc.git 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 libgomp/ * testsuite/libgomp.oacc-c-c++-common/atomic_capture-2.c: Remove iteration-ordering assumptions. --- 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 842f2de4722..4f83f03899d 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; }