[gcc(refs/vendors/ARM/heads/arm-perf-staging)] [OpenACC 'exit data'] Evaluate 'copyfrom' individually for 'GOMP_MAP_STRUCT' entries
Tamar Christina
tnfchris@gcc.gnu.org
Fri Jul 17 15:21:15 GMT 2020
https://gcc.gnu.org/g:4664ca1bc40318dbe60591cfe6d31c3d36d439c3
commit 4664ca1bc40318dbe60591cfe6d31c3d36d439c3
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Thu Jun 4 16:13:35 2020 +0200
[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)
Diff:
---
libgomp/oacc-mem.c | 16 ++++
.../testsuite/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 a34f4cf0e91..11419e692aa 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 285be84f244..543aaa15306 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
}
More information about the Gcc-cvs
mailing list