From 2c838a3e4ea06c69c856d074ae5b0400e08ae3c2 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge 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) : 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 #include @@ -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