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

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index b7c85cf5976f2d2de5aa145bbd476416810a2c7e..a34f4cf0e91860cadeb4255e9d48d266ea355078 100644 (file)
--- 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 (file)
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 (file)
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;
-}