Hi!

On 2019-12-17T22:03:47-0800, Julian Brown <jul...@codesourcery.com> wrote:
> This part contains the libgomp runtime support for the GOMP_MAP_ATTACH and
> GOMP_MAP_DETACH mapping kinds (etc.)

> --- a/libgomp/oacc-mem.c
> +++ b/libgomp/oacc-mem.c

> @@ -1075,6 +1119,39 @@ goacc_exit_data_internal (struct gomp_device_descr 
> *acc_dev, size_t mapnum,

> +     case GOMP_MAP_STRUCT:
> +       {
> +         int elems = sizes[i];
> +         for (int j = 1; j <= elems; j++)
> +           {
> +             struct splay_tree_key_s k;
> +             k.host_start = (uintptr_t) hostaddrs[i + j];
> +             k.host_end = k.host_start + sizes[i + j];
> +             splay_tree_key str;
> +             str = splay_tree_lookup (&acc_dev->mem_map, &k);
> +             if (str)
> +               {
> +                 if (finalize)
> +                   {
> +                     str->refcount -= str->virtual_refcount;
> +                     str->virtual_refcount = 0;
> +                   }
> +                 if (str->virtual_refcount > 0)
> +                   {
> +                     str->refcount--;
> +                     str->virtual_refcount--;
> +                   }
> +                 else if (str->refcount > 0)
> +                   str->refcount--;
> +                 if (str->refcount == 0)
> +                   gomp_remove_var_async (acc_dev, str, aq);
> +               }
> +           }
> +         i += elems;
> +       }
> +       break;

I'm aware that this 'GOMP_MAP_STRUCT' special handling shouldn't have
been there to begin with, and is now scheduled to go away (yay!), but
while testing a few things while reviewing (reverse-engineering the
intentions of) these fix-up patches, I quickly ran into cases where
OpenACC code that I understand to be valid failed, exactly here.  I've
pushed "[OpenACC 'exit data'] Evaluate 'finalize' individually for
'GOMP_MAP_STRUCT' entries" to master branch in commit
a02f1adbfe619ab19cf142438e0a02950d3594da, and releases/gcc-10 branch in
commit 5a1b479aedd83d0362f870f480a24a011e703de4, and then "[OpenACC 'exit
data'] Evaluate 'copyfrom' individually for 'GOMP_MAP_STRUCT' entries" to
master branch in commit 2c838a3e4ea06c69c856d074ae5b0400e08ae3c2, and
releases/gcc-10 branch in commit
4664ca1bc40318dbe60591cfe6d31c3d36d439c3, see attached.


Grüße
 Thomas


-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander 
Walter
>From a02f1adbfe619ab19cf142438e0a02950d3594da Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Thu, 4 Jun 2020 16:01:07 +0200
Subject: [PATCH] [OpenACC 'exit data'] Evaluate 'finalize' individually for
 'GOMP_MAP_STRUCT' entries

Currently, we don't at all evaluate 'finalize' for 'GOMP_MAP_STRUCT' entries.
Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code.

	libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>:
	Evaluate 'finalize' individually for each entry.
	* testsuite/libgomp.oacc-c-c++-common/struct-1.c: New file.
	* testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: Remove
	file.
---
 libgomp/oacc-mem.c                            |  10 ++
 .../libgomp.oacc-c-c++-common/struct-1.c      | 146 ++++++++++++++++++
 .../struct-refcount-1.c                       |  47 ------
 3 files changed, 156 insertions(+), 47 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
 delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index b7c85cf5976f..a34f4cf0e918 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1184,6 +1184,16 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	    int elems = sizes[i];
 	    for (int j = 1; j <= elems; j++)
 	      {
+		assert (i + j < mapnum);
+
+		kind = kinds[i + j] & 0xff;
+
+		finalize = false;
+		if (kind == GOMP_MAP_FORCE_FROM
+		    || kind == GOMP_MAP_DELETE
+		    || kind == GOMP_MAP_FORCE_DETACH)
+		  finalize = true;
+
 		struct splay_tree_key_s k;
 		k.host_start = (uintptr_t) hostaddrs[i + j];
 		k.host_end = k.host_start + sizes[i + j];
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
new file mode 100644
index 000000000000..285be84f244b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
@@ -0,0 +1,146 @@
+/* Test dynamic refcount of separate structure members.  */
+
+#include <assert.h>
+#include <stdbool.h>
+#include <openacc.h>
+
+struct s
+{
+  signed char a;
+  float b;
+};
+
+static void test(unsigned variant)
+{
+  struct s s;
+
+#pragma acc enter data create(s.a, s.b)
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+
+  if (variant & 4)
+    {
+      if (variant & 8)
+	{
+#pragma acc enter data create(s.b)
+	}
+      else
+	acc_create(&s.b, sizeof s.b);
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+
+      if (variant & 16)
+	{
+#pragma acc enter data create(s.a)
+	}
+      else
+	acc_create(&s.a, sizeof s.a);
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+
+      if (variant & 32)
+	{
+#pragma acc enter data create(s.a)
+	  acc_create(&s.b, sizeof s.b);
+#pragma acc enter data create(s.b)
+#pragma acc enter data create(s.b)
+	  acc_create(&s.a, sizeof s.a);
+	  acc_create(&s.a, sizeof s.a);
+	  acc_create(&s.a, sizeof s.a);
+	}
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+    }
+
+#pragma acc parallel \
+  copy(s.a, s.b)
+  {
+  }
+
+  if (variant & 32)
+    {
+      if (variant & 1)
+	{
+#pragma acc exit data delete(s.a) finalize
+	}
+      else
+	acc_delete_finalize(&s.a, sizeof s.a);
+    }
+  else
+    {
+      if (variant & 1)
+	{
+#pragma acc exit data delete(s.a)
+	}
+      else
+	acc_delete(&s.a, sizeof s.a);
+      if (variant & 4)
+	{
+	  assert(acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+	  if (variant & 1)
+	    {
+#pragma acc exit data delete(s.a)
+	    }
+	  else
+	    acc_delete(&s.a, sizeof s.a);
+	}
+    }
+#if ACC_MEM_SHARED
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+#else
+  assert(!acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+#endif
+
+  if (variant & 32)
+    {
+      if (variant & 2)
+	{
+#pragma acc exit data delete(s.b) finalize
+	}
+      else
+	acc_delete_finalize(&s.b, sizeof s.b);
+    }
+  else
+    {
+      if (variant & 2)
+	{
+#pragma acc exit data delete(s.b)
+	}
+      else
+	acc_delete(&s.b, sizeof s.b);
+      if (variant & 4)
+	{
+#if ACC_MEM_SHARED
+	  assert(acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+#else
+	  assert(!acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+#endif
+	  if (variant & 2)
+	    {
+#pragma acc exit data delete(s.b)
+	    }
+	  else
+	    acc_delete(&s.b, sizeof s.b);
+	}
+    }
+#if ACC_MEM_SHARED
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+#else
+  assert(!acc_is_present(&s.a, sizeof s.a));
+  assert(!acc_is_present(&s.b, sizeof s.b));
+#endif
+}
+
+int main()
+{
+  for (unsigned variant = 0; variant < 64; ++variant)
+    test(variant);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
deleted file mode 100644
index bde5890d6676..000000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
+++ /dev/null
@@ -1,47 +0,0 @@
-/* Test dynamic unmapping of separate structure members.  */
-
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-
-#include <assert.h>
-#include <stdbool.h>
-#include <openacc.h>
-
-struct s
-{
-  char a;
-  float b;
-};
-
-void test (bool use_directives)
-{
-  struct s s;
-
-#pragma acc enter data create(s.a, s.b)
-  assert (acc_is_present (&s.a, sizeof s.a));
-  assert (acc_is_present (&s.b, sizeof s.b));
-
-  if (use_directives)
-    {
-#pragma acc exit data delete(s.a)
-    }
-  else
-    acc_delete (&s.a, sizeof s.a);
-  assert (!acc_is_present (&s.a, sizeof s.a));
-  assert (acc_is_present (&s.b, sizeof s.b));
-  if (use_directives)
-    {
-#pragma acc exit data delete(s.b)
-    }
-  else
-    acc_delete (&s.b, sizeof s.b);
-  assert (!acc_is_present (&s.a, sizeof s.a));
-  assert (!acc_is_present (&s.b, sizeof s.b));
-}
-
-int main ()
-{
-  test (true);
-  test (false);
-
-  return 0;
-}
-- 
2.26.2

>From 5a1b479aedd83d0362f870f480a24a011e703de4 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Thu, 4 Jun 2020 16:01:07 +0200
Subject: [PATCH] [OpenACC 'exit data'] Evaluate 'finalize' individually for
 'GOMP_MAP_STRUCT' entries

Currently, we don't at all evaluate 'finalize' for 'GOMP_MAP_STRUCT' entries.
Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code.

	libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>:
	Evaluate 'finalize' individually for each entry.
	* testsuite/libgomp.oacc-c-c++-common/struct-1.c: New file.
	* testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c: Remove
	file.

(cherry picked from commit a02f1adbfe619ab19cf142438e0a02950d3594da)
---
 libgomp/oacc-mem.c                            |  10 ++
 .../libgomp.oacc-c-c++-common/struct-1.c      | 146 ++++++++++++++++++
 .../struct-refcount-1.c                       |  47 ------
 3 files changed, 156 insertions(+), 47 deletions(-)
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
 delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index b7c85cf5976f..a34f4cf0e918 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1184,6 +1184,16 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 	    int elems = sizes[i];
 	    for (int j = 1; j <= elems; j++)
 	      {
+		assert (i + j < mapnum);
+
+		kind = kinds[i + j] & 0xff;
+
+		finalize = false;
+		if (kind == GOMP_MAP_FORCE_FROM
+		    || kind == GOMP_MAP_DELETE
+		    || kind == GOMP_MAP_FORCE_DETACH)
+		  finalize = true;
+
 		struct splay_tree_key_s k;
 		k.host_start = (uintptr_t) hostaddrs[i + j];
 		k.host_end = k.host_start + sizes[i + j];
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
new file mode 100644
index 000000000000..285be84f244b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
@@ -0,0 +1,146 @@
+/* Test dynamic refcount of separate structure members.  */
+
+#include <assert.h>
+#include <stdbool.h>
+#include <openacc.h>
+
+struct s
+{
+  signed char a;
+  float b;
+};
+
+static void test(unsigned variant)
+{
+  struct s s;
+
+#pragma acc enter data create(s.a, s.b)
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+
+  if (variant & 4)
+    {
+      if (variant & 8)
+	{
+#pragma acc enter data create(s.b)
+	}
+      else
+	acc_create(&s.b, sizeof s.b);
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+
+      if (variant & 16)
+	{
+#pragma acc enter data create(s.a)
+	}
+      else
+	acc_create(&s.a, sizeof s.a);
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+
+      if (variant & 32)
+	{
+#pragma acc enter data create(s.a)
+	  acc_create(&s.b, sizeof s.b);
+#pragma acc enter data create(s.b)
+#pragma acc enter data create(s.b)
+	  acc_create(&s.a, sizeof s.a);
+	  acc_create(&s.a, sizeof s.a);
+	  acc_create(&s.a, sizeof s.a);
+	}
+      assert(acc_is_present(&s.a, sizeof s.a));
+      assert(acc_is_present(&s.b, sizeof s.b));
+    }
+
+#pragma acc parallel \
+  copy(s.a, s.b)
+  {
+  }
+
+  if (variant & 32)
+    {
+      if (variant & 1)
+	{
+#pragma acc exit data delete(s.a) finalize
+	}
+      else
+	acc_delete_finalize(&s.a, sizeof s.a);
+    }
+  else
+    {
+      if (variant & 1)
+	{
+#pragma acc exit data delete(s.a)
+	}
+      else
+	acc_delete(&s.a, sizeof s.a);
+      if (variant & 4)
+	{
+	  assert(acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+	  if (variant & 1)
+	    {
+#pragma acc exit data delete(s.a)
+	    }
+	  else
+	    acc_delete(&s.a, sizeof s.a);
+	}
+    }
+#if ACC_MEM_SHARED
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+#else
+  assert(!acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+#endif
+
+  if (variant & 32)
+    {
+      if (variant & 2)
+	{
+#pragma acc exit data delete(s.b) finalize
+	}
+      else
+	acc_delete_finalize(&s.b, sizeof s.b);
+    }
+  else
+    {
+      if (variant & 2)
+	{
+#pragma acc exit data delete(s.b)
+	}
+      else
+	acc_delete(&s.b, sizeof s.b);
+      if (variant & 4)
+	{
+#if ACC_MEM_SHARED
+	  assert(acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+#else
+	  assert(!acc_is_present(&s.a, sizeof s.a));
+	  assert(acc_is_present(&s.b, sizeof s.b));
+#endif
+	  if (variant & 2)
+	    {
+#pragma acc exit data delete(s.b)
+	    }
+	  else
+	    acc_delete(&s.b, sizeof s.b);
+	}
+    }
+#if ACC_MEM_SHARED
+  assert(acc_is_present(&s.a, sizeof s.a));
+  assert(acc_is_present(&s.b, sizeof s.b));
+#else
+  assert(!acc_is_present(&s.a, sizeof s.a));
+  assert(!acc_is_present(&s.b, sizeof s.b));
+#endif
+}
+
+int main()
+{
+  for (unsigned variant = 0; variant < 64; ++variant)
+    test(variant);
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
deleted file mode 100644
index bde5890d6676..000000000000
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c
+++ /dev/null
@@ -1,47 +0,0 @@
-/* Test dynamic unmapping of separate structure members.  */
-
-/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
-
-#include <assert.h>
-#include <stdbool.h>
-#include <openacc.h>
-
-struct s
-{
-  char a;
-  float b;
-};
-
-void test (bool use_directives)
-{
-  struct s s;
-
-#pragma acc enter data create(s.a, s.b)
-  assert (acc_is_present (&s.a, sizeof s.a));
-  assert (acc_is_present (&s.b, sizeof s.b));
-
-  if (use_directives)
-    {
-#pragma acc exit data delete(s.a)
-    }
-  else
-    acc_delete (&s.a, sizeof s.a);
-  assert (!acc_is_present (&s.a, sizeof s.a));
-  assert (acc_is_present (&s.b, sizeof s.b));
-  if (use_directives)
-    {
-#pragma acc exit data delete(s.b)
-    }
-  else
-    acc_delete (&s.b, sizeof s.b);
-  assert (!acc_is_present (&s.a, sizeof s.a));
-  assert (!acc_is_present (&s.b, sizeof s.b));
-}
-
-int main ()
-{
-  test (true);
-  test (false);
-
-  return 0;
-}
-- 
2.26.2

>From 2c838a3e4ea06c69c856d074ae5b0400e08ae3c2 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Thu, 4 Jun 2020 16:13:35 +0200
Subject: [PATCH] [OpenACC 'exit data'] Evaluate 'copyfrom' individually for
 'GOMP_MAP_STRUCT' entries

Currently, we don't at all evaluate 'copyfrom' for 'GOMP_MAP_STRUCT' entries.
Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code.

	libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>:
	Evaluate 'copyfrom' individually for each entry.
	* testsuite/libgomp.oacc-c-c++-common/struct-1.c: Update.
---
 libgomp/oacc-mem.c                            | 16 ++++
 .../libgomp.oacc-c-c++-common/struct-1.c      | 93 +++++++++++++------
 2 files changed, 83 insertions(+), 26 deletions(-)

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index a34f4cf0e918..11419e692aa2 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1194,6 +1194,12 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 		    || kind == GOMP_MAP_FORCE_DETACH)
 		  finalize = true;
 
+		copyfrom = false;
+		if (kind == GOMP_MAP_FROM
+		    || kind == GOMP_MAP_FORCE_FROM
+		    || kind == GOMP_MAP_ALWAYS_FROM)
+		  copyfrom = true;
+
 		struct splay_tree_key_s k;
 		k.host_start = (uintptr_t) hostaddrs[i + j];
 		k.host_end = k.host_start + sizes[i + j];
@@ -1216,6 +1222,16 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 		    else if (str->refcount > 0
 			     && str->refcount != REFCOUNT_INFINITY)
 		      str->refcount--;
+
+		    if (copyfrom
+			&& (kind != GOMP_MAP_FROM || str->refcount == 0))
+		      gomp_copy_dev2host (acc_dev, aq, (void *) k.host_start,
+					  (void *) (str->tgt->tgt_start
+						    + str->tgt_offset
+						    + k.host_start
+						    - str->host_start),
+					  k.host_end - k.host_start);
+
 		    if (str->refcount == 0)
 		      {
 			if (aq)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
index 285be84f244b..543aaa153064 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
@@ -1,4 +1,4 @@
-/* Test dynamic refcount of separate structure members.  */
+/* Test dynamic refcount and copy behavior of separate structure members.  */
 
 #include <assert.h>
 #include <stdbool.h>
@@ -12,41 +12,45 @@ struct s
 
 static void test(unsigned variant)
 {
-  struct s s;
+  struct s s = { .a = 73, .b = -22 };
 
-#pragma acc enter data create(s.a, s.b)
+#pragma acc enter data copyin(s.a, s.b)
   assert(acc_is_present(&s.a, sizeof s.a));
   assert(acc_is_present(&s.b, sizeof s.b));
 
+  /* To verify that any following 'copyin' doesn't 'copyin' again.  */
+  s.a = -s.a;
+  s.b = -s.b;
+
   if (variant & 4)
     {
       if (variant & 8)
 	{
-#pragma acc enter data create(s.b)
+#pragma acc enter data copyin(s.b)
 	}
       else
-	acc_create(&s.b, sizeof s.b);
+	acc_copyin(&s.b, sizeof s.b);
       assert(acc_is_present(&s.a, sizeof s.a));
       assert(acc_is_present(&s.b, sizeof s.b));
 
       if (variant & 16)
 	{
-#pragma acc enter data create(s.a)
+#pragma acc enter data copyin(s.a)
 	}
       else
-	acc_create(&s.a, sizeof s.a);
+	acc_copyin(&s.a, sizeof s.a);
       assert(acc_is_present(&s.a, sizeof s.a));
       assert(acc_is_present(&s.b, sizeof s.b));
 
       if (variant & 32)
 	{
-#pragma acc enter data create(s.a)
-	  acc_create(&s.b, sizeof s.b);
-#pragma acc enter data create(s.b)
-#pragma acc enter data create(s.b)
-	  acc_create(&s.a, sizeof s.a);
-	  acc_create(&s.a, sizeof s.a);
-	  acc_create(&s.a, sizeof s.a);
+#pragma acc enter data copyin(s.a)
+	  acc_copyin(&s.b, sizeof s.b);
+#pragma acc enter data copyin(s.b)
+#pragma acc enter data copyin(s.b)
+	  acc_copyin(&s.a, sizeof s.a);
+	  acc_copyin(&s.a, sizeof s.a);
+	  acc_copyin(&s.a, sizeof s.a);
 	}
       assert(acc_is_present(&s.a, sizeof s.a));
       assert(acc_is_present(&s.b, sizeof s.b));
@@ -55,85 +59,122 @@ static void test(unsigned variant)
 #pragma acc parallel \
   copy(s.a, s.b)
   {
+#if ACC_MEM_SHARED
+    if (s.a++ != -73)
+      __builtin_abort();
+    if (s.b-- != 22)
+      __builtin_abort();
+#else
+    if (s.a++ != 73)
+      __builtin_abort();
+    if (s.b-- != -22)
+      __builtin_abort();
+#endif
   }
+#if ACC_MEM_SHARED
+  assert(s.a == -72);
+  assert(s.b == 21);
+#else
+  assert(s.a == -73);
+  assert(s.b == 22);
+#endif
 
   if (variant & 32)
     {
       if (variant & 1)
 	{
-#pragma acc exit data delete(s.a) finalize
+#pragma acc exit data copyout(s.a) finalize
 	}
       else
-	acc_delete_finalize(&s.a, sizeof s.a);
+	acc_copyout_finalize(&s.a, sizeof s.a);
     }
   else
     {
       if (variant & 1)
 	{
-#pragma acc exit data delete(s.a)
+#pragma acc exit data copyout(s.a)
 	}
       else
-	acc_delete(&s.a, sizeof s.a);
+	acc_copyout(&s.a, sizeof s.a);
       if (variant & 4)
 	{
 	  assert(acc_is_present(&s.a, sizeof s.a));
 	  assert(acc_is_present(&s.b, sizeof s.b));
+#if ACC_MEM_SHARED
+	  assert(s.a == -72);
+	  assert(s.b == 21);
+#else
+	  assert(s.a == -73);
+	  assert(s.b == 22);
+#endif
 	  if (variant & 1)
 	    {
-#pragma acc exit data delete(s.a)
+#pragma acc exit data copyout(s.a)
 	    }
 	  else
-	    acc_delete(&s.a, sizeof s.a);
+	    acc_copyout(&s.a, sizeof s.a);
 	}
     }
 #if ACC_MEM_SHARED
   assert(acc_is_present(&s.a, sizeof s.a));
   assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == -72);
+  assert(s.b == 21);
 #else
   assert(!acc_is_present(&s.a, sizeof s.a));
   assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == 74);
+  assert(s.b == 22);
 #endif
 
   if (variant & 32)
     {
       if (variant & 2)
 	{
-#pragma acc exit data delete(s.b) finalize
+#pragma acc exit data copyout(s.b) finalize
 	}
       else
-	acc_delete_finalize(&s.b, sizeof s.b);
+	acc_copyout_finalize(&s.b, sizeof s.b);
     }
   else
     {
       if (variant & 2)
 	{
-#pragma acc exit data delete(s.b)
+#pragma acc exit data copyout(s.b)
 	}
       else
-	acc_delete(&s.b, sizeof s.b);
+	acc_copyout(&s.b, sizeof s.b);
       if (variant & 4)
 	{
 #if ACC_MEM_SHARED
 	  assert(acc_is_present(&s.a, sizeof s.a));
 	  assert(acc_is_present(&s.b, sizeof s.b));
+	  assert(s.a == -72);
+	  assert(s.b == 21);
 #else
 	  assert(!acc_is_present(&s.a, sizeof s.a));
 	  assert(acc_is_present(&s.b, sizeof s.b));
+	  assert(s.a == 74);
+	  assert(s.b == 22);
 #endif
 	  if (variant & 2)
 	    {
-#pragma acc exit data delete(s.b)
+#pragma acc exit data copyout(s.b)
 	    }
 	  else
-	    acc_delete(&s.b, sizeof s.b);
+	    acc_copyout(&s.b, sizeof s.b);
 	}
     }
 #if ACC_MEM_SHARED
   assert(acc_is_present(&s.a, sizeof s.a));
   assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == -72);
+  assert(s.b == 21);
 #else
   assert(!acc_is_present(&s.a, sizeof s.a));
   assert(!acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == 74);
+  assert(s.b == -23);
 #endif
 }
 
-- 
2.26.2

>From 4664ca1bc40318dbe60591cfe6d31c3d36d439c3 Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <tho...@codesourcery.com>
Date: Thu, 4 Jun 2020 16:13:35 +0200
Subject: [PATCH] [OpenACC 'exit data'] Evaluate 'copyfrom' individually for
 'GOMP_MAP_STRUCT' entries

Currently, we don't at all evaluate 'copyfrom' for 'GOMP_MAP_STRUCT' entries.
Fix this by copying/adapting the corresponding non-'GOMP_MAP_STRUCT' code.

	libgomp/
	* oacc-mem.c (goacc_exit_data_internal) <GOMP_MAP_STRUCT>:
	Evaluate 'copyfrom' individually for each entry.
	* testsuite/libgomp.oacc-c-c++-common/struct-1.c: Update.

(cherry picked from commit 2c838a3e4ea06c69c856d074ae5b0400e08ae3c2)
---
 libgomp/oacc-mem.c                            | 16 ++++
 .../libgomp.oacc-c-c++-common/struct-1.c      | 93 +++++++++++++------
 2 files changed, 83 insertions(+), 26 deletions(-)

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index a34f4cf0e918..11419e692aa2 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1194,6 +1194,12 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 		    || kind == GOMP_MAP_FORCE_DETACH)
 		  finalize = true;
 
+		copyfrom = false;
+		if (kind == GOMP_MAP_FROM
+		    || kind == GOMP_MAP_FORCE_FROM
+		    || kind == GOMP_MAP_ALWAYS_FROM)
+		  copyfrom = true;
+
 		struct splay_tree_key_s k;
 		k.host_start = (uintptr_t) hostaddrs[i + j];
 		k.host_end = k.host_start + sizes[i + j];
@@ -1216,6 +1222,16 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum,
 		    else if (str->refcount > 0
 			     && str->refcount != REFCOUNT_INFINITY)
 		      str->refcount--;
+
+		    if (copyfrom
+			&& (kind != GOMP_MAP_FROM || str->refcount == 0))
+		      gomp_copy_dev2host (acc_dev, aq, (void *) k.host_start,
+					  (void *) (str->tgt->tgt_start
+						    + str->tgt_offset
+						    + k.host_start
+						    - str->host_start),
+					  k.host_end - k.host_start);
+
 		    if (str->refcount == 0)
 		      {
 			if (aq)
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
index 285be84f244b..543aaa153064 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c
@@ -1,4 +1,4 @@
-/* Test dynamic refcount of separate structure members.  */
+/* Test dynamic refcount and copy behavior of separate structure members.  */
 
 #include <assert.h>
 #include <stdbool.h>
@@ -12,41 +12,45 @@ struct s
 
 static void test(unsigned variant)
 {
-  struct s s;
+  struct s s = { .a = 73, .b = -22 };
 
-#pragma acc enter data create(s.a, s.b)
+#pragma acc enter data copyin(s.a, s.b)
   assert(acc_is_present(&s.a, sizeof s.a));
   assert(acc_is_present(&s.b, sizeof s.b));
 
+  /* To verify that any following 'copyin' doesn't 'copyin' again.  */
+  s.a = -s.a;
+  s.b = -s.b;
+
   if (variant & 4)
     {
       if (variant & 8)
 	{
-#pragma acc enter data create(s.b)
+#pragma acc enter data copyin(s.b)
 	}
       else
-	acc_create(&s.b, sizeof s.b);
+	acc_copyin(&s.b, sizeof s.b);
       assert(acc_is_present(&s.a, sizeof s.a));
       assert(acc_is_present(&s.b, sizeof s.b));
 
       if (variant & 16)
 	{
-#pragma acc enter data create(s.a)
+#pragma acc enter data copyin(s.a)
 	}
       else
-	acc_create(&s.a, sizeof s.a);
+	acc_copyin(&s.a, sizeof s.a);
       assert(acc_is_present(&s.a, sizeof s.a));
       assert(acc_is_present(&s.b, sizeof s.b));
 
       if (variant & 32)
 	{
-#pragma acc enter data create(s.a)
-	  acc_create(&s.b, sizeof s.b);
-#pragma acc enter data create(s.b)
-#pragma acc enter data create(s.b)
-	  acc_create(&s.a, sizeof s.a);
-	  acc_create(&s.a, sizeof s.a);
-	  acc_create(&s.a, sizeof s.a);
+#pragma acc enter data copyin(s.a)
+	  acc_copyin(&s.b, sizeof s.b);
+#pragma acc enter data copyin(s.b)
+#pragma acc enter data copyin(s.b)
+	  acc_copyin(&s.a, sizeof s.a);
+	  acc_copyin(&s.a, sizeof s.a);
+	  acc_copyin(&s.a, sizeof s.a);
 	}
       assert(acc_is_present(&s.a, sizeof s.a));
       assert(acc_is_present(&s.b, sizeof s.b));
@@ -55,85 +59,122 @@ static void test(unsigned variant)
 #pragma acc parallel \
   copy(s.a, s.b)
   {
+#if ACC_MEM_SHARED
+    if (s.a++ != -73)
+      __builtin_abort();
+    if (s.b-- != 22)
+      __builtin_abort();
+#else
+    if (s.a++ != 73)
+      __builtin_abort();
+    if (s.b-- != -22)
+      __builtin_abort();
+#endif
   }
+#if ACC_MEM_SHARED
+  assert(s.a == -72);
+  assert(s.b == 21);
+#else
+  assert(s.a == -73);
+  assert(s.b == 22);
+#endif
 
   if (variant & 32)
     {
       if (variant & 1)
 	{
-#pragma acc exit data delete(s.a) finalize
+#pragma acc exit data copyout(s.a) finalize
 	}
       else
-	acc_delete_finalize(&s.a, sizeof s.a);
+	acc_copyout_finalize(&s.a, sizeof s.a);
     }
   else
     {
       if (variant & 1)
 	{
-#pragma acc exit data delete(s.a)
+#pragma acc exit data copyout(s.a)
 	}
       else
-	acc_delete(&s.a, sizeof s.a);
+	acc_copyout(&s.a, sizeof s.a);
       if (variant & 4)
 	{
 	  assert(acc_is_present(&s.a, sizeof s.a));
 	  assert(acc_is_present(&s.b, sizeof s.b));
+#if ACC_MEM_SHARED
+	  assert(s.a == -72);
+	  assert(s.b == 21);
+#else
+	  assert(s.a == -73);
+	  assert(s.b == 22);
+#endif
 	  if (variant & 1)
 	    {
-#pragma acc exit data delete(s.a)
+#pragma acc exit data copyout(s.a)
 	    }
 	  else
-	    acc_delete(&s.a, sizeof s.a);
+	    acc_copyout(&s.a, sizeof s.a);
 	}
     }
 #if ACC_MEM_SHARED
   assert(acc_is_present(&s.a, sizeof s.a));
   assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == -72);
+  assert(s.b == 21);
 #else
   assert(!acc_is_present(&s.a, sizeof s.a));
   assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == 74);
+  assert(s.b == 22);
 #endif
 
   if (variant & 32)
     {
       if (variant & 2)
 	{
-#pragma acc exit data delete(s.b) finalize
+#pragma acc exit data copyout(s.b) finalize
 	}
       else
-	acc_delete_finalize(&s.b, sizeof s.b);
+	acc_copyout_finalize(&s.b, sizeof s.b);
     }
   else
     {
       if (variant & 2)
 	{
-#pragma acc exit data delete(s.b)
+#pragma acc exit data copyout(s.b)
 	}
       else
-	acc_delete(&s.b, sizeof s.b);
+	acc_copyout(&s.b, sizeof s.b);
       if (variant & 4)
 	{
 #if ACC_MEM_SHARED
 	  assert(acc_is_present(&s.a, sizeof s.a));
 	  assert(acc_is_present(&s.b, sizeof s.b));
+	  assert(s.a == -72);
+	  assert(s.b == 21);
 #else
 	  assert(!acc_is_present(&s.a, sizeof s.a));
 	  assert(acc_is_present(&s.b, sizeof s.b));
+	  assert(s.a == 74);
+	  assert(s.b == 22);
 #endif
 	  if (variant & 2)
 	    {
-#pragma acc exit data delete(s.b)
+#pragma acc exit data copyout(s.b)
 	    }
 	  else
-	    acc_delete(&s.b, sizeof s.b);
+	    acc_copyout(&s.b, sizeof s.b);
 	}
     }
 #if ACC_MEM_SHARED
   assert(acc_is_present(&s.a, sizeof s.a));
   assert(acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == -72);
+  assert(s.b == 21);
 #else
   assert(!acc_is_present(&s.a, sizeof s.a));
   assert(!acc_is_present(&s.b, sizeof s.b));
+  assert(s.a == 74);
+  assert(s.b == -23);
 #endif
 }
 
-- 
2.26.2

Reply via email to