Hi! On 2019-12-17T22:03:47-0800, Julian Brown <jul...@codesourcery.com> wrote: > This part contains the libgomp runtime support for the GOMP_MAP_ATTACH and > GOMP_MAP_DETACH mapping kinds (etc.)
> --- a/libgomp/oacc-mem.c > +++ b/libgomp/oacc-mem.c > @@ -1075,6 +1119,39 @@ goacc_exit_data_internal (struct gomp_device_descr > *acc_dev, size_t mapnum, > + case GOMP_MAP_STRUCT: > + { > + int elems = sizes[i]; > + for (int j = 1; j <= elems; j++) > + { > + struct splay_tree_key_s k; > + k.host_start = (uintptr_t) hostaddrs[i + j]; > + k.host_end = k.host_start + sizes[i + j]; > + splay_tree_key str; > + str = splay_tree_lookup (&acc_dev->mem_map, &k); > + if (str) > + { > + if (finalize) > + { > + str->refcount -= str->virtual_refcount; > + str->virtual_refcount = 0; > + } > + if (str->virtual_refcount > 0) > + { > + str->refcount--; > + str->virtual_refcount--; > + } > + else if (str->refcount > 0) > + str->refcount--; > + if (str->refcount == 0) > + gomp_remove_var_async (acc_dev, str, aq); > + } > + } > + i += elems; > + } > + break; I'm aware that this 'GOMP_MAP_STRUCT' special handling shouldn't have been there to begin with, and is now scheduled to go away (yay!), but while testing a few things while reviewing (reverse-engineering the intentions of) these fix-up patches, I quickly ran into cases where OpenACC code that I understand to be valid failed, exactly here. I've pushed "[OpenACC 'exit data'] Evaluate 'finalize' individually for 'GOMP_MAP_STRUCT' entries" to master branch in commit a02f1adbfe619ab19cf142438e0a02950d3594da, and releases/gcc-10 branch in commit 5a1b479aedd83d0362f870f480a24a011e703de4, and then "[OpenACC 'exit data'] Evaluate 'copyfrom' individually for 'GOMP_MAP_STRUCT' entries" to master branch in commit 2c838a3e4ea06c69c856d074ae5b0400e08ae3c2, and releases/gcc-10 branch in commit 4664ca1bc40318dbe60591cfe6d31c3d36d439c3, see attached. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
>From a02f1adbfe619ab19cf142438e0a02950d3594da Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <tho...@codesourcery.com> Date: Thu, 4 Jun 2020 16:01:07 +0200 Subject: [PATCH] [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. --- libgomp/oacc-mem.c | 10 ++ .../libgomp.oacc-c-c++-common/struct-1.c | 146 ++++++++++++++++++ .../struct-refcount-1.c | 47 ------ 3 files changed, 156 insertions(+), 47 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index b7c85cf5976f..a34f4cf0e918 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 000000000000..285be84f244b --- /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 bde5890d6676..000000000000 --- 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; -} -- 2.26.2
>From 5a1b479aedd83d0362f870f480a24a011e703de4 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <tho...@codesourcery.com> Date: Thu, 4 Jun 2020 16:01:07 +0200 Subject: [PATCH] [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. (cherry picked from commit a02f1adbfe619ab19cf142438e0a02950d3594da) --- libgomp/oacc-mem.c | 10 ++ .../libgomp.oacc-c-c++-common/struct-1.c | 146 ++++++++++++++++++ .../struct-refcount-1.c | 47 ------ 3 files changed, 156 insertions(+), 47 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-1.c delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/struct-refcount-1.c diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index b7c85cf5976f..a34f4cf0e918 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 000000000000..285be84f244b --- /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 bde5890d6676..000000000000 --- 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; -} -- 2.26.2
>From 2c838a3e4ea06c69c856d074ae5b0400e08ae3c2 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <tho...@codesourcery.com> 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) <GOMP_MAP_STRUCT>: 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 <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 } -- 2.26.2
>From 4664ca1bc40318dbe60591cfe6d31c3d36d439c3 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <tho...@codesourcery.com> 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) <GOMP_MAP_STRUCT>: Evaluate 'copyfrom' individually for each entry. * testsuite/libgomp.oacc-c-c++-common/struct-1.c: Update. (cherry picked from commit 2c838a3e4ea06c69c856d074ae5b0400e08ae3c2) --- 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 <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 } -- 2.26.2