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 --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 b976094998f2df8765e15c1c964f6c46d67c4c49..b8a76a17560d08565407888a7c4ba2d9171d07d1 100644 (file)
--- 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; }