https://gcc.gnu.org/g:2b033568b996ec934055310c8626f71bd6992138
commit r16-5573-g2b033568b996ec934055310c8626f71bd6992138 Author: Arsen Arsenović <[email protected]> Date: Mon Nov 24 13:35:18 2025 +0100 libgomp/oacc: fix atomic_capture-3 iteration ordering issues 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. Diff: --- .../libgomp.oacc-c-c++-common/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; }
