PR92036 "OpenACC 'firstprivate' clause: initial value"

PR92036 "OpenACC 'firstprivate' clause: initial value"

[Thomas Schwinge]

[-- Attachment #1.1: Type: text/plain, Size: 271 bytes --]

Hi!

In r276757 "[PR92036] Add
'libgomp.oacc-c-c++-common/data-firstprivate-1.c'", I committed the
attached to document the status quo, and then what I think it instead
should be per PR92036 "OpenACC 'firstprivate' clause: initial value".


Grüße
 Thomas



[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #1.2: 0001-PR92036-Add-libgomp.oacc-c-c-common-data-first.trunk.patch --]
[-- Type: text/x-diff, Size: 5397 bytes --]

From 5fc105387fda327bf5e36ae6f997e415da6d1f37 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed, 9 Oct 2019 11:31:14 +0000
Subject: [PATCH] [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.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@276757 138bc75d-0d04-0410-961f-82ee72b054a4
---
 libgomp/ChangeLog                             |   6 +
 .../data-firstprivate-1.c                     | 165 ++++++++++++++++++
 2 files changed, 171 insertions(+)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/data-firstprivate-1.c

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 1b43c456741..319a1911882 100644
--- 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
index 00000000000..8900a4e070d
--- /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;
+}
-- 
2.17.1


[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 658 bytes --]