[gcc r11-3059] openacc: Fix atomic_capture-2.c iteration-ordering issues

Julian Brown jules@gcc.gnu.org
Tue Sep 8 20:28:00 GMT 2020


https://gcc.gnu.org/g:8183ebcdc1c843f15c807e5bc26dbafe4e8c4dc3

commit r11-3059-g8183ebcdc1c843f15c807e5bc26dbafe4e8c4dc3
Author: Julian Brown <julian@codesourcery.com>
Date:   Mon Sep 7 11:43:16 2020 -0700

    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.

Diff:
---
 .../libgomp.oacc-c-c++-common/atomic_capture-2.c   | 92 ++++++++++------------
 1 file 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 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; }


More information about the Gcc-cvs mailing list