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

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index a34f4cf0e91860cadeb4255e9d48d266ea355078..11419e692aa2e7f0967bec4f1eff70bca6d69665 100644 (file)
--- 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 285be84f244b1925aee29459153b714f01d9cb87..543aaa153064683ec22d10ef02b0bfa562edba5f 100644 (file)
--- 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
 }