[openacc, testsuite, PR85527, committed] Fix undefined behaviour in atomic_capture-1.c

[Tom de Vries]

Hi,

atomic_capture-1.c contains the following bit:
...
  fgot = 1.0;
  fexp = 0.0;

#pragma acc data copy (fgot, fdata[0:N])
  {
#pragma acc parallel loop
    for (i = 0; i < N; i++)
      {
        float expr = 32.0;

#pragma acc atomic capture
        fdata[i] = fgot = expr - fgot;
      }
  }

  for (i = 0; i < N; i++)
    if (i % 2 == 0)
      {
        if (fdata[i] != 31.0)
          abort ();
      }
    else
      {
        if (fdata[i] != 1.0)
          abort ();
      }
...

The problem here is that the checking code assumes that the atomic 
operations execute in a certain order, while that is not guaranteed by 
the openacc standard.

This check happens to pass for nvptx accelerator, but not for gcn 
accelerator.

Fixed by Julian's patch below, which makes the check order-independent.

Tested on x86_64 with nvptx accelerator.

Committed to trunk.

Thanks,
- Tom
[openacc, testsuite] Fix undefined behaviour in atomic_capture-1.c

2018-04-29  Julian Brown  <jul...@codesourcery.com>
	    Tom de Vries  <t...@codesourcery.com>

	PR testsuite/85527
	* testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c: Allow
	arbitrary order for iterations of atomic subtract check.

---
 .../libgomp.oacc-c-c++-common/atomic_capture-1.c      | 19 +++++++++----------
 1 file changed, 9 insertions(+), 10 deletions(-)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c
index ad958cd..9b71a08 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/atomic_capture-1.c
@@ -795,17 +795,16 @@ main(int argc, char **argv)
       }
   }
 
+  int ones = 0, thirtyones = 0;
+
   for (i = 0; i < N; i++)
-    if (i % 2 == 0)
-      {
-	if (fdata[i] != 31.0)
-	  abort ();
-      }
-    else
-      {
-	if (fdata[i] != 1.0)
-	  abort ();
-      }
+    if (fdata[i] == 31.0)
+      thirtyones++;
+    else if (fdata[i] == 1.0)
+      ones++;
+
+  if (ones != N / 2 || thirtyones != N / 2)
+    abort ();
 
 
   /* BINOP = / */