[gcc r12-908] Add 'libgomp.oacc-c-c++-common/private-atomic-1.c' [PR83812]

[gcc r12-908] Add 'libgomp.oacc-c-c++-common/private-atomic-1.c' [PR83812]

[Thomas Schwinge]

https://gcc.gnu.org/g:1467100fc72562a59f70cdd4e05f6c810d1fadcc

commit r12-908-g1467100fc72562a59f70cdd4e05f6c810d1fadcc
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Wed May 19 11:58:49 2021 +0200

    Add 'libgomp.oacc-c-c++-common/private-atomic-1.c' [PR83812]
    
    ... to at least document/test/XFAIL nvptx offloading: PR83812 "operation not
    supported on global/shared address space".
    
            libgomp/
            PR target/83812
            * testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c: New.

Diff:
---
 .../libgomp.oacc-c-c++-common/private-atomic-1.c   | 37 ++++++++++++++++++++++
 1 file changed, 37 insertions(+)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c
new file mode 100644
index 00000000000..77197d8fd44
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/private-atomic-1.c
@@ -0,0 +1,37 @@
+// 'atomic' access of thread-private variable
+
+#include <assert.h>
+
+int main (void)
+{
+  int res;
+
+  res = 0;
+#pragma acc parallel reduction(+: res)
+  {
+#pragma acc loop vector reduction(+: res)
+    for (int i = 0; i < 2322; i++)
+    {
+      int v = -222;
+
+#pragma acc loop seq
+      for (int j = 0; j < 121; ++j)
+	{
+#pragma acc atomic update
+	  ++v;
+	  /* nvptx offloading: PR83812 "operation not supported on global/shared address space".
+	     { dg-output "(\n|\r\n|\r)libgomp: cuStreamSynchronize error: operation not supported on global/shared address space(\n|\r\n|\r)$" { target openacc_nvidia_accel_selected } }
+	       Scan for what we expect in the "XFAILed" case (without actually XFAILing).
+	     { dg-shouldfail "XFAILed" { openacc_nvidia_accel_selected } }
+	       ... instead of 'dg-xfail-run-if' so that 'dg-output' is evaluated at all.
+	     { dg-final { if { [dg-process-target { xfail openacc_nvidia_accel_selected }] == "F" } { xfail "[testname-for-summary] really is XFAILed" } } }
+	       ... so that we still get an XFAIL visible in the log.  */
+	}
+
+      res += (v == -222 + 121);
+    }
+  }
+  assert (res == 2322);
+
+  return 0;
+}