[PR92036] Add 'libgomp.oacc-c-c++-common/data-firstprivate-1.c'

	libgomp/
	PR middle-end/92036
	* testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New
	file.

From-SVN: r276757

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 1b43c456741e26396fe51a9bfb1ca9b681f82798..319a19118825298f17f429238b658fc7d202f31c 100644 (file)
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,9 @@
+2019-10-09  Thomas Schwinge  <thomas@codesourcery.com>
+
+       PR middle-end/92036
+       * testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c: New
+       file.
+
 2019-10-09  Tobias Burnus  <tobias@codesourcery.com>
 
        PR testsuite/91884

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c
new file mode 100644 (file)
index 0000000..8900a4e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c
@@ -0,0 +1,165 @@
+/* Test behavior of 'firstprivate' lexically vs. dynamically nested inside a
+   'data' region.  */
+
+#include <stdlib.h>
+
+
+#define VERIFY(x) \
+  do { \
+    if (!(x)) \
+      abort (); \
+  } while (0);
+
+
+/* This is basically and extended version of 't2' from 'firstprivate-1.c'.  */
+
+int lexically_nested_val = 2;
+
+static void
+lexically_nested ()
+{
+#pragma acc data \
+  copy (lexically_nested_val)
+  {
+    VERIFY (lexically_nested_val == 2);
+
+#pragma acc parallel \
+  present (lexically_nested_val)
+    {
+      VERIFY (lexically_nested_val == 2);
+
+      /* This updates the device copy, or shared variable.  */
+      lexically_nested_val = 7;
+    }
+
+#if ACC_MEM_SHARED
+    VERIFY (lexically_nested_val == 7);
+#else
+    VERIFY (lexically_nested_val == 2);
+#endif
+
+    /* This only updates the local/shared variable, but not the device
+       copy.  */
+    lexically_nested_val = 5;
+
+#pragma acc parallel \
+    firstprivate (lexically_nested_val)
+      {
+#if 1 /* Current behavior.  */
+       /* The 'firstprivate' copy is initialized from the device copy, or
+          shared variable.  */
+# if ACC_MEM_SHARED
+       VERIFY (lexically_nested_val == 5);
+# else
+       VERIFY (lexically_nested_val == 7);
+# endif
+#else /* Expected behavior per PR92036.  */
+       /* The 'firstprivate' copy is initialized from the local thread.  */
+       VERIFY (lexically_nested_val == 5);
+#endif
+
+       /* This updates the 'firstprivate' copy only, but not the shared
+          variable.  */
+       lexically_nested_val = 9;
+      }
+
+    VERIFY (lexically_nested_val == 5);
+  }
+  /* If not shared, the device copy has now been copied back.  */
+
+#if ACC_MEM_SHARED
+  VERIFY (lexically_nested_val == 5);
+#else
+  VERIFY (lexically_nested_val == 7);
+#endif
+}
+
+
+int dynamically_nested_val = 2;
+
+/* Same as above, but compute construct 1 broken out, so no longer lexically
+   nested inside 'data' region.  */
+
+static void
+dynamically_nested_compute_1 ()
+{
+#pragma acc parallel \
+  present (dynamically_nested_val)
+  {
+    VERIFY (dynamically_nested_val == 2);
+
+    /* This updates the device copy, or shared variable.  */
+    dynamically_nested_val = 7;
+  }
+}
+
+/* Same as above, but compute construct 2 broken out, so no longer lexically
+   nested inside 'data' region.  */
+
+static void
+dynamically_nested_compute_2 ()
+{
+#pragma acc parallel \
+  firstprivate (dynamically_nested_val)
+    {
+#if 1 /* Current behavior.  */
+      /* The 'firstprivate' copy is initialized from the device copy, or shared
+        variable.  */
+# if ACC_MEM_SHARED
+      VERIFY (dynamically_nested_val == 5);
+# else
+      VERIFY (dynamically_nested_val == 7);
+# endif
+#else /* Expected behavior per PR92036.  */
+      /* The 'firstprivate' copy is initialized from the local thread.  */
+      VERIFY (dynamically_nested_val == 5);
+#endif
+
+      /* This updates the 'firstprivate' copy only, but not the shared
+        variable.  */
+      dynamically_nested_val = 9;
+    }
+}
+
+static void
+dynamically_nested ()
+{
+#pragma acc data \
+  copy (dynamically_nested_val)
+  {
+    VERIFY (dynamically_nested_val == 2);
+
+    dynamically_nested_compute_1 ();
+
+#if ACC_MEM_SHARED
+    VERIFY (dynamically_nested_val == 7);
+#else
+    VERIFY (dynamically_nested_val == 2);
+#endif
+
+    /* This only updates the local/shared variable, but not the device
+       copy.  */
+    dynamically_nested_val = 5;
+
+    dynamically_nested_compute_2 ();
+
+    VERIFY (dynamically_nested_val == 5);
+  }
+  /* If not shared, the device copy has now been copied back.  */
+
+#if ACC_MEM_SHARED
+  VERIFY (dynamically_nested_val == 5);
+#else
+  VERIFY (dynamically_nested_val == 7);
+#endif
+}
+
+
+int
+main()
+{
+  lexically_nested ();
+  dynamically_nested ();
+
+  return 0;
+}