[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.

3 files changed, 156 insertions, 47 deletions

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index b7c85cf..a34f4cf 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 0000000..285be84
--- /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 bde5890..0000000
--- 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;
-}